LLVM  9.0.0svn
AMDGPUHSAMetadataStreamer.cpp
Go to the documentation of this file.
1 //===--- AMDGPUHSAMetadataStreamer.cpp --------------------------*- C++ -*-===//
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 /// \file
10 /// AMDGPU HSA Metadata Streamer.
11 ///
12 //
13 //===----------------------------------------------------------------------===//
14 
16 #include "AMDGPU.h"
17 #include "AMDGPUSubtarget.h"
19 #include "SIMachineFunctionInfo.h"
20 #include "SIProgramInfo.h"
21 #include "Utils/AMDGPUBaseInfo.h"
22 #include "llvm/ADT/StringSwitch.h"
23 #include "llvm/IR/Constants.h"
24 #include "llvm/IR/Module.h"
26 
27 namespace llvm {
28 
29 static cl::opt<bool> DumpHSAMetadata(
30  "amdgpu-dump-hsa-metadata",
31  cl::desc("Dump AMDGPU HSA Metadata"));
32 static cl::opt<bool> VerifyHSAMetadata(
33  "amdgpu-verify-hsa-metadata",
34  cl::desc("Verify AMDGPU HSA Metadata"));
35 
36 namespace AMDGPU {
37 namespace HSAMD {
38 
39 //===----------------------------------------------------------------------===//
40 // HSAMetadataStreamerV2
41 //===----------------------------------------------------------------------===//
42 void MetadataStreamerV2::dump(StringRef HSAMetadataString) const {
43  errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
44 }
45 
46 void MetadataStreamerV2::verify(StringRef HSAMetadataString) const {
47  errs() << "AMDGPU HSA Metadata Parser Test: ";
48 
49  HSAMD::Metadata FromHSAMetadataString;
50  if (fromString(HSAMetadataString, FromHSAMetadataString)) {
51  errs() << "FAIL\n";
52  return;
53  }
54 
55  std::string ToHSAMetadataString;
56  if (toString(FromHSAMetadataString, ToHSAMetadataString)) {
57  errs() << "FAIL\n";
58  return;
59  }
60 
61  errs() << (HSAMetadataString == ToHSAMetadataString ? "PASS" : "FAIL")
62  << '\n';
63  if (HSAMetadataString != ToHSAMetadataString) {
64  errs() << "Original input: " << HSAMetadataString << '\n'
65  << "Produced output: " << ToHSAMetadataString << '\n';
66  }
67 }
68 
70 MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const {
71  if (AccQual.empty())
73 
74  return StringSwitch<AccessQualifier>(AccQual)
75  .Case("read_only", AccessQualifier::ReadOnly)
76  .Case("write_only", AccessQualifier::WriteOnly)
77  .Case("read_write", AccessQualifier::ReadWrite)
78  .Default(AccessQualifier::Default);
79 }
80 
82 MetadataStreamerV2::getAddressSpaceQualifier(
83  unsigned AddressSpace) const {
84  switch (AddressSpace) {
97  default:
99  }
100 }
101 
102 ValueKind MetadataStreamerV2::getValueKind(Type *Ty, StringRef TypeQual,
103  StringRef BaseTypeName) const {
104  if (TypeQual.find("pipe") != StringRef::npos)
105  return ValueKind::Pipe;
106 
107  return StringSwitch<ValueKind>(BaseTypeName)
108  .Case("image1d_t", ValueKind::Image)
109  .Case("image1d_array_t", ValueKind::Image)
110  .Case("image1d_buffer_t", ValueKind::Image)
111  .Case("image2d_t", ValueKind::Image)
112  .Case("image2d_array_t", ValueKind::Image)
113  .Case("image2d_array_depth_t", ValueKind::Image)
114  .Case("image2d_array_msaa_t", ValueKind::Image)
115  .Case("image2d_array_msaa_depth_t", ValueKind::Image)
116  .Case("image2d_depth_t", ValueKind::Image)
117  .Case("image2d_msaa_t", ValueKind::Image)
118  .Case("image2d_msaa_depth_t", ValueKind::Image)
119  .Case("image3d_t", ValueKind::Image)
120  .Case("sampler_t", ValueKind::Sampler)
121  .Case("queue_t", ValueKind::Queue)
122  .Default(isa<PointerType>(Ty) ?
123  (Ty->getPointerAddressSpace() ==
128 }
129 
130 ValueType MetadataStreamerV2::getValueType(Type *Ty, StringRef TypeName) const {
131  switch (Ty->getTypeID()) {
132  case Type::IntegerTyID: {
133  auto Signed = !TypeName.startswith("u");
134  switch (Ty->getIntegerBitWidth()) {
135  case 8:
137  case 16:
139  case 32:
141  case 64:
143  default:
144  return ValueType::Struct;
145  }
146  }
147  case Type::HalfTyID:
148  return ValueType::F16;
149  case Type::FloatTyID:
150  return ValueType::F32;
151  case Type::DoubleTyID:
152  return ValueType::F64;
153  case Type::PointerTyID:
154  return getValueType(Ty->getPointerElementType(), TypeName);
155  case Type::VectorTyID:
156  return getValueType(Ty->getVectorElementType(), TypeName);
157  default:
158  return ValueType::Struct;
159  }
160 }
161 
162 std::string MetadataStreamerV2::getTypeName(Type *Ty, bool Signed) const {
163  switch (Ty->getTypeID()) {
164  case Type::IntegerTyID: {
165  if (!Signed)
166  return (Twine('u') + getTypeName(Ty, true)).str();
167 
168  auto BitWidth = Ty->getIntegerBitWidth();
169  switch (BitWidth) {
170  case 8:
171  return "char";
172  case 16:
173  return "short";
174  case 32:
175  return "int";
176  case 64:
177  return "long";
178  default:
179  return (Twine('i') + Twine(BitWidth)).str();
180  }
181  }
182  case Type::HalfTyID:
183  return "half";
184  case Type::FloatTyID:
185  return "float";
186  case Type::DoubleTyID:
187  return "double";
188  case Type::VectorTyID: {
189  auto VecTy = cast<VectorType>(Ty);
190  auto ElTy = VecTy->getElementType();
191  auto NumElements = VecTy->getVectorNumElements();
192  return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
193  }
194  default:
195  return "unknown";
196  }
197 }
198 
199 std::vector<uint32_t>
200 MetadataStreamerV2::getWorkGroupDimensions(MDNode *Node) const {
201  std::vector<uint32_t> Dims;
202  if (Node->getNumOperands() != 3)
203  return Dims;
204 
205  for (auto &Op : Node->operands())
206  Dims.push_back(mdconst::extract<ConstantInt>(Op)->getZExtValue());
207  return Dims;
208 }
209 
210 Kernel::CodeProps::Metadata
211 MetadataStreamerV2::getHSACodeProps(const MachineFunction &MF,
212  const SIProgramInfo &ProgramInfo) const {
213  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
214  const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
215  HSAMD::Kernel::CodeProps::Metadata HSACodeProps;
216  const Function &F = MF.getFunction();
217 
218  assert(F.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
219  F.getCallingConv() == CallingConv::SPIR_KERNEL);
220 
221  unsigned MaxKernArgAlign;
222  HSACodeProps.mKernargSegmentSize = STM.getKernArgSegmentSize(F,
223  MaxKernArgAlign);
224  HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize;
225  HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize;
226  HSACodeProps.mKernargSegmentAlign = std::max(MaxKernArgAlign, 4u);
227  HSACodeProps.mWavefrontSize = STM.getWavefrontSize();
228  HSACodeProps.mNumSGPRs = ProgramInfo.NumSGPR;
229  HSACodeProps.mNumVGPRs = ProgramInfo.NumVGPR;
230  HSACodeProps.mMaxFlatWorkGroupSize = MFI.getMaxFlatWorkGroupSize();
231  HSACodeProps.mIsDynamicCallStack = ProgramInfo.DynamicCallStack;
232  HSACodeProps.mIsXNACKEnabled = STM.isXNACKEnabled();
233  HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs();
234  HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs();
235 
236  return HSACodeProps;
237 }
238 
239 Kernel::DebugProps::Metadata
240 MetadataStreamerV2::getHSADebugProps(const MachineFunction &MF,
241  const SIProgramInfo &ProgramInfo) const {
242  return HSAMD::Kernel::DebugProps::Metadata();
243 }
244 
245 void MetadataStreamerV2::emitVersion() {
246  auto &Version = HSAMetadata.mVersion;
247 
248  Version.push_back(VersionMajor);
249  Version.push_back(VersionMinor);
250 }
251 
252 void MetadataStreamerV2::emitPrintf(const Module &Mod) {
253  auto &Printf = HSAMetadata.mPrintf;
254 
255  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
256  if (!Node)
257  return;
258 
259  for (auto Op : Node->operands())
260  if (Op->getNumOperands())
261  Printf.push_back(cast<MDString>(Op->getOperand(0))->getString());
262 }
263 
264 void MetadataStreamerV2::emitKernelLanguage(const Function &Func) {
265  auto &Kernel = HSAMetadata.mKernels.back();
266 
267  // TODO: What about other languages?
268  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
269  if (!Node || !Node->getNumOperands())
270  return;
271  auto Op0 = Node->getOperand(0);
272  if (Op0->getNumOperands() <= 1)
273  return;
274 
275  Kernel.mLanguage = "OpenCL C";
276  Kernel.mLanguageVersion.push_back(
277  mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue());
278  Kernel.mLanguageVersion.push_back(
279  mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue());
280 }
281 
282 void MetadataStreamerV2::emitKernelAttrs(const Function &Func) {
283  auto &Attrs = HSAMetadata.mKernels.back().mAttrs;
284 
285  if (auto Node = Func.getMetadata("reqd_work_group_size"))
286  Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node);
287  if (auto Node = Func.getMetadata("work_group_size_hint"))
288  Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node);
289  if (auto Node = Func.getMetadata("vec_type_hint")) {
290  Attrs.mVecTypeHint = getTypeName(
291  cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
292  mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue());
293  }
294  if (Func.hasFnAttribute("runtime-handle")) {
295  Attrs.mRuntimeHandle =
296  Func.getFnAttribute("runtime-handle").getValueAsString().str();
297  }
298 }
299 
300 void MetadataStreamerV2::emitKernelArgs(const Function &Func) {
301  for (auto &Arg : Func.args())
302  emitKernelArg(Arg);
303 
304  emitHiddenKernelArgs(Func);
305 }
306 
307 void MetadataStreamerV2::emitKernelArg(const Argument &Arg) {
308  auto Func = Arg.getParent();
309  auto ArgNo = Arg.getArgNo();
310  const MDNode *Node;
311 
312  StringRef Name;
313  Node = Func->getMetadata("kernel_arg_name");
314  if (Node && ArgNo < Node->getNumOperands())
315  Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
316  else if (Arg.hasName())
317  Name = Arg.getName();
318 
319  StringRef TypeName;
320  Node = Func->getMetadata("kernel_arg_type");
321  if (Node && ArgNo < Node->getNumOperands())
322  TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
323 
324  StringRef BaseTypeName;
325  Node = Func->getMetadata("kernel_arg_base_type");
326  if (Node && ArgNo < Node->getNumOperands())
327  BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
328 
329  StringRef AccQual;
330  if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
331  Arg.hasNoAliasAttr()) {
332  AccQual = "read_only";
333  } else {
334  Node = Func->getMetadata("kernel_arg_access_qual");
335  if (Node && ArgNo < Node->getNumOperands())
336  AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
337  }
338 
339  StringRef TypeQual;
340  Node = Func->getMetadata("kernel_arg_type_qual");
341  if (Node && ArgNo < Node->getNumOperands())
342  TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
343 
344  Type *Ty = Arg.getType();
345  const DataLayout &DL = Func->getParent()->getDataLayout();
346 
347  unsigned PointeeAlign = 0;
348  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
349  if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
350  PointeeAlign = Arg.getParamAlignment();
351  if (PointeeAlign == 0)
352  PointeeAlign = DL.getABITypeAlignment(PtrTy->getElementType());
353  }
354  }
355 
356  emitKernelArg(DL, Ty, getValueKind(Arg.getType(), TypeQual, BaseTypeName),
357  PointeeAlign, Name, TypeName, BaseTypeName, AccQual, TypeQual);
358 }
359 
360 void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty,
362  unsigned PointeeAlign, StringRef Name,
363  StringRef TypeName,
364  StringRef BaseTypeName,
365  StringRef AccQual, StringRef TypeQual) {
366  HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
367  auto &Arg = HSAMetadata.mKernels.back().mArgs.back();
368 
369  Arg.mName = Name;
370  Arg.mTypeName = TypeName;
371  Arg.mSize = DL.getTypeAllocSize(Ty);
372  Arg.mAlign = DL.getABITypeAlignment(Ty);
373  Arg.mValueKind = ValueKind;
374  Arg.mValueType = getValueType(Ty, BaseTypeName);
375  Arg.mPointeeAlign = PointeeAlign;
376 
377  if (auto PtrTy = dyn_cast<PointerType>(Ty))
378  Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace());
379 
380  Arg.mAccQual = getAccessQualifier(AccQual);
381 
382  // TODO: Emit Arg.mActualAccQual.
383 
384  SmallVector<StringRef, 1> SplitTypeQuals;
385  TypeQual.split(SplitTypeQuals, " ", -1, false);
386  for (StringRef Key : SplitTypeQuals) {
387  auto P = StringSwitch<bool*>(Key)
388  .Case("const", &Arg.mIsConst)
389  .Case("restrict", &Arg.mIsRestrict)
390  .Case("volatile", &Arg.mIsVolatile)
391  .Case("pipe", &Arg.mIsPipe)
392  .Default(nullptr);
393  if (P)
394  *P = true;
395  }
396 }
397 
398 void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func) {
399  int HiddenArgNumBytes =
400  getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
401 
402  if (!HiddenArgNumBytes)
403  return;
404 
405  auto &DL = Func.getParent()->getDataLayout();
406  auto Int64Ty = Type::getInt64Ty(Func.getContext());
407 
408  if (HiddenArgNumBytes >= 8)
409  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetX);
410  if (HiddenArgNumBytes >= 16)
411  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetY);
412  if (HiddenArgNumBytes >= 24)
413  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetZ);
414 
415  auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
417 
418  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
419  // "none" argument.
420  if (HiddenArgNumBytes >= 32) {
421  if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
422  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenPrintfBuffer);
423  else
424  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
425  }
426 
427  // Emit "default queue" and "completion action" arguments if enqueue kernel is
428  // used, otherwise emit dummy "none" arguments.
429  if (HiddenArgNumBytes >= 48) {
430  if (Func.hasFnAttribute("calls-enqueue-kernel")) {
431  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenDefaultQueue);
432  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenCompletionAction);
433  } else {
434  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
435  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
436  }
437  }
438 }
439 
441  return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
442 }
443 
445  emitVersion();
446  emitPrintf(Mod);
447 }
448 
450  std::string HSAMetadataString;
451  if (toString(HSAMetadata, HSAMetadataString))
452  return;
453 
454  if (DumpHSAMetadata)
455  dump(HSAMetadataString);
456  if (VerifyHSAMetadata)
457  verify(HSAMetadataString);
458 }
459 
461  const SIProgramInfo &ProgramInfo) {
462  auto &Func = MF.getFunction();
463  if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL)
464  return;
465 
466  auto CodeProps = getHSACodeProps(MF, ProgramInfo);
467  auto DebugProps = getHSADebugProps(MF, ProgramInfo);
468 
469  HSAMetadata.mKernels.push_back(Kernel::Metadata());
470  auto &Kernel = HSAMetadata.mKernels.back();
471 
472  Kernel.mName = Func.getName();
473  Kernel.mSymbolName = (Twine(Func.getName()) + Twine("@kd")).str();
474  emitKernelLanguage(Func);
475  emitKernelAttrs(Func);
476  emitKernelArgs(Func);
477  HSAMetadata.mKernels.back().mCodeProps = CodeProps;
478  HSAMetadata.mKernels.back().mDebugProps = DebugProps;
479 }
480 
481 //===----------------------------------------------------------------------===//
482 // HSAMetadataStreamerV3
483 //===----------------------------------------------------------------------===//
484 
485 void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
486  errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
487 }
488 
489 void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
490  errs() << "AMDGPU HSA Metadata Parser Test: ";
491 
492  msgpack::Document FromHSAMetadataString;
493 
494  if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
495  errs() << "FAIL\n";
496  return;
497  }
498 
499  std::string ToHSAMetadataString;
500  raw_string_ostream StrOS(ToHSAMetadataString);
501  FromHSAMetadataString.toYAML(StrOS);
502 
503  errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
504  if (HSAMetadataString != ToHSAMetadataString) {
505  errs() << "Original input: " << HSAMetadataString << '\n'
506  << "Produced output: " << StrOS.str() << '\n';
507  }
508 }
509 
511 MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const {
513  .Case("read_only", StringRef("read_only"))
514  .Case("write_only", StringRef("write_only"))
515  .Case("read_write", StringRef("read_write"))
516  .Default(None);
517 }
518 
520 MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const {
521  switch (AddressSpace) {
523  return StringRef("private");
525  return StringRef("global");
527  return StringRef("constant");
529  return StringRef("local");
531  return StringRef("generic");
533  return StringRef("region");
534  default:
535  return None;
536  }
537 }
538 
539 StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual,
540  StringRef BaseTypeName) const {
541  if (TypeQual.find("pipe") != StringRef::npos)
542  return "pipe";
543 
544  return StringSwitch<StringRef>(BaseTypeName)
545  .Case("image1d_t", "image")
546  .Case("image1d_array_t", "image")
547  .Case("image1d_buffer_t", "image")
548  .Case("image2d_t", "image")
549  .Case("image2d_array_t", "image")
550  .Case("image2d_array_depth_t", "image")
551  .Case("image2d_array_msaa_t", "image")
552  .Case("image2d_array_msaa_depth_t", "image")
553  .Case("image2d_depth_t", "image")
554  .Case("image2d_msaa_t", "image")
555  .Case("image2d_msaa_depth_t", "image")
556  .Case("image3d_t", "image")
557  .Case("sampler_t", "sampler")
558  .Case("queue_t", "queue")
559  .Default(isa<PointerType>(Ty)
561  ? "dynamic_shared_pointer"
562  : "global_buffer")
563  : "by_value");
564 }
565 
566 StringRef MetadataStreamerV3::getValueType(Type *Ty, StringRef TypeName) const {
567  switch (Ty->getTypeID()) {
568  case Type::IntegerTyID: {
569  auto Signed = !TypeName.startswith("u");
570  switch (Ty->getIntegerBitWidth()) {
571  case 8:
572  return Signed ? "i8" : "u8";
573  case 16:
574  return Signed ? "i16" : "u16";
575  case 32:
576  return Signed ? "i32" : "u32";
577  case 64:
578  return Signed ? "i64" : "u64";
579  default:
580  return "struct";
581  }
582  }
583  case Type::HalfTyID:
584  return "f16";
585  case Type::FloatTyID:
586  return "f32";
587  case Type::DoubleTyID:
588  return "f64";
589  case Type::PointerTyID:
590  return getValueType(Ty->getPointerElementType(), TypeName);
591  case Type::VectorTyID:
592  return getValueType(Ty->getVectorElementType(), TypeName);
593  default:
594  return "struct";
595  }
596 }
597 
598 std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
599  switch (Ty->getTypeID()) {
600  case Type::IntegerTyID: {
601  if (!Signed)
602  return (Twine('u') + getTypeName(Ty, true)).str();
603 
604  auto BitWidth = Ty->getIntegerBitWidth();
605  switch (BitWidth) {
606  case 8:
607  return "char";
608  case 16:
609  return "short";
610  case 32:
611  return "int";
612  case 64:
613  return "long";
614  default:
615  return (Twine('i') + Twine(BitWidth)).str();
616  }
617  }
618  case Type::HalfTyID:
619  return "half";
620  case Type::FloatTyID:
621  return "float";
622  case Type::DoubleTyID:
623  return "double";
624  case Type::VectorTyID: {
625  auto VecTy = cast<VectorType>(Ty);
626  auto ElTy = VecTy->getElementType();
627  auto NumElements = VecTy->getVectorNumElements();
628  return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
629  }
630  default:
631  return "unknown";
632  }
633 }
634 
636 MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
637  auto Dims = HSAMetadataDoc->getArrayNode();
638  if (Node->getNumOperands() != 3)
639  return Dims;
640 
641  for (auto &Op : Node->operands())
642  Dims.push_back(Dims.getDocument()->getNode(
643  uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
644  return Dims;
645 }
646 
647 void MetadataStreamerV3::emitVersion() {
648  auto Version = HSAMetadataDoc->getArrayNode();
649  Version.push_back(Version.getDocument()->getNode(VersionMajor));
650  Version.push_back(Version.getDocument()->getNode(VersionMinor));
651  getRootMetadata("amdhsa.version") = Version;
652 }
653 
654 void MetadataStreamerV3::emitPrintf(const Module &Mod) {
655  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
656  if (!Node)
657  return;
658 
659  auto Printf = HSAMetadataDoc->getArrayNode();
660  for (auto Op : Node->operands())
661  if (Op->getNumOperands())
662  Printf.push_back(Printf.getDocument()->getNode(
663  cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
664  getRootMetadata("amdhsa.printf") = Printf;
665 }
666 
667 void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
668  msgpack::MapDocNode Kern) {
669  // TODO: What about other languages?
670  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
671  if (!Node || !Node->getNumOperands())
672  return;
673  auto Op0 = Node->getOperand(0);
674  if (Op0->getNumOperands() <= 1)
675  return;
676 
677  Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
678  auto LanguageVersion = Kern.getDocument()->getArrayNode();
679  LanguageVersion.push_back(Kern.getDocument()->getNode(
680  mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
681  LanguageVersion.push_back(Kern.getDocument()->getNode(
682  mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
683  Kern[".language_version"] = LanguageVersion;
684 }
685 
686 void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
687  msgpack::MapDocNode Kern) {
688 
689  if (auto Node = Func.getMetadata("reqd_work_group_size"))
690  Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
691  if (auto Node = Func.getMetadata("work_group_size_hint"))
692  Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
693  if (auto Node = Func.getMetadata("vec_type_hint")) {
694  Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
695  getTypeName(
696  cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
697  mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
698  /*Copy=*/true);
699  }
700  if (Func.hasFnAttribute("runtime-handle")) {
701  Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
702  Func.getFnAttribute("runtime-handle").getValueAsString().str(),
703  /*Copy=*/true);
704  }
705 }
706 
707 void MetadataStreamerV3::emitKernelArgs(const Function &Func,
708  msgpack::MapDocNode Kern) {
709  unsigned Offset = 0;
710  auto Args = HSAMetadataDoc->getArrayNode();
711  for (auto &Arg : Func.args())
712  emitKernelArg(Arg, Offset, Args);
713 
714  emitHiddenKernelArgs(Func, Offset, Args);
715 
716  Kern[".args"] = Args;
717 }
718 
719 void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
721  auto Func = Arg.getParent();
722  auto ArgNo = Arg.getArgNo();
723  const MDNode *Node;
724 
725  StringRef Name;
726  Node = Func->getMetadata("kernel_arg_name");
727  if (Node && ArgNo < Node->getNumOperands())
728  Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
729  else if (Arg.hasName())
730  Name = Arg.getName();
731 
733  Node = Func->getMetadata("kernel_arg_type");
734  if (Node && ArgNo < Node->getNumOperands())
735  TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
736 
737  StringRef BaseTypeName;
738  Node = Func->getMetadata("kernel_arg_base_type");
739  if (Node && ArgNo < Node->getNumOperands())
740  BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
741 
743  if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
744  Arg.hasNoAliasAttr()) {
745  AccQual = "read_only";
746  } else {
747  Node = Func->getMetadata("kernel_arg_access_qual");
748  if (Node && ArgNo < Node->getNumOperands())
749  AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
750  }
751 
752  StringRef TypeQual;
753  Node = Func->getMetadata("kernel_arg_type_qual");
754  if (Node && ArgNo < Node->getNumOperands())
755  TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
756 
757  Type *Ty = Arg.getType();
758  const DataLayout &DL = Func->getParent()->getDataLayout();
759 
760  unsigned PointeeAlign = 0;
761  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
762  if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
763  PointeeAlign = Arg.getParamAlignment();
764  if (PointeeAlign == 0)
765  PointeeAlign = DL.getABITypeAlignment(PtrTy->getElementType());
766  }
767  }
768 
769  emitKernelArg(Func->getParent()->getDataLayout(), Arg.getType(),
770  getValueKind(Arg.getType(), TypeQual, BaseTypeName), Offset,
771  Args, PointeeAlign, Name, TypeName, BaseTypeName, AccQual,
772  TypeQual);
773 }
774 
775 void MetadataStreamerV3::emitKernelArg(const DataLayout &DL, Type *Ty,
776  StringRef ValueKind, unsigned &Offset,
778  unsigned PointeeAlign, StringRef Name,
779  StringRef TypeName,
780  StringRef BaseTypeName,
781  StringRef AccQual, StringRef TypeQual) {
782  auto Arg = Args.getDocument()->getMapNode();
783 
784  if (!Name.empty())
785  Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
786  if (!TypeName.empty())
787  Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
788  auto Size = DL.getTypeAllocSize(Ty);
789  auto Align = DL.getABITypeAlignment(Ty);
790  Arg[".size"] = Arg.getDocument()->getNode(Size);
791  Offset = alignTo(Offset, Align);
792  Arg[".offset"] = Arg.getDocument()->getNode(Offset);
793  Offset += Size;
794  Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
795  Arg[".value_type"] =
796  Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true);
797  if (PointeeAlign)
798  Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign);
799 
800  if (auto PtrTy = dyn_cast<PointerType>(Ty))
801  if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
802  Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
803 
804  if (auto AQ = getAccessQualifier(AccQual))
805  Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
806 
807  // TODO: Emit Arg[".actual_access"].
808 
809  SmallVector<StringRef, 1> SplitTypeQuals;
810  TypeQual.split(SplitTypeQuals, " ", -1, false);
811  for (StringRef Key : SplitTypeQuals) {
812  if (Key == "const")
813  Arg[".is_const"] = Arg.getDocument()->getNode(true);
814  else if (Key == "restrict")
815  Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
816  else if (Key == "volatile")
817  Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
818  else if (Key == "pipe")
819  Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
820  }
821 
822  Args.push_back(Arg);
823 }
824 
825 void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
826  unsigned &Offset,
827  msgpack::ArrayDocNode Args) {
828  int HiddenArgNumBytes =
829  getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
830 
831  if (!HiddenArgNumBytes)
832  return;
833 
834  auto &DL = Func.getParent()->getDataLayout();
835  auto Int64Ty = Type::getInt64Ty(Func.getContext());
836 
837  if (HiddenArgNumBytes >= 8)
838  emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, Args);
839  if (HiddenArgNumBytes >= 16)
840  emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, Args);
841  if (HiddenArgNumBytes >= 24)
842  emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, Args);
843 
844  auto Int8PtrTy =
846 
847  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
848  // "none" argument.
849  if (HiddenArgNumBytes >= 32) {
850  if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
851  emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, Args);
852  else
853  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
854  }
855 
856  // Emit "default queue" and "completion action" arguments if enqueue kernel is
857  // used, otherwise emit dummy "none" arguments.
858  if (HiddenArgNumBytes >= 48) {
859  if (Func.hasFnAttribute("calls-enqueue-kernel")) {
860  emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args);
861  emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
862  } else {
863  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
864  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
865  }
866  }
867 }
868 
870 MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
871  const SIProgramInfo &ProgramInfo) const {
872  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
874  const Function &F = MF.getFunction();
875 
876  auto Kern = HSAMetadataDoc->getMapNode();
877 
878  unsigned MaxKernArgAlign;
879  Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
880  STM.getKernArgSegmentSize(F, MaxKernArgAlign));
881  Kern[".group_segment_fixed_size"] =
882  Kern.getDocument()->getNode(ProgramInfo.LDSSize);
883  Kern[".private_segment_fixed_size"] =
884  Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
885  Kern[".kernarg_segment_align"] =
886  Kern.getDocument()->getNode(std::max(uint32_t(4), MaxKernArgAlign));
887  Kern[".wavefront_size"] =
888  Kern.getDocument()->getNode(STM.getWavefrontSize());
889  Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
890  Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
891  Kern[".max_flat_workgroup_size"] =
892  Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
893  Kern[".sgpr_spill_count"] =
894  Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
895  Kern[".vgpr_spill_count"] =
896  Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
897 
898  return Kern;
899 }
900 
902  return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
903 }
904 
906  emitVersion();
907  emitPrintf(Mod);
908  getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
909 }
910 
912  std::string HSAMetadataString;
913  raw_string_ostream StrOS(HSAMetadataString);
914  HSAMetadataDoc->toYAML(StrOS);
915 
916  if (DumpHSAMetadata)
917  dump(StrOS.str());
918  if (VerifyHSAMetadata)
919  verify(StrOS.str());
920 }
921 
923  const SIProgramInfo &ProgramInfo) {
924  auto &Func = MF.getFunction();
925  auto Kern = getHSAKernelProps(MF, ProgramInfo);
926 
929 
930  auto Kernels =
931  getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
932 
933  {
934  Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
935  Kern[".symbol"] = Kern.getDocument()->getNode(
936  (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
937  emitKernelLanguage(Func, Kern);
938  emitKernelAttrs(Func, Kern);
939  emitKernelArgs(Func, Kern);
940  }
941 
942  Kernels.push_back(Kern);
943 }
944 
945 } // end namespace HSAMD
946 } // end namespace AMDGPU
947 } // end namespace llvm
Type * getVectorElementType() const
Definition: Type.h:370
const NoneType None
Definition: None.h:23
Address space for global memory (RAT0, VTX0).
Definition: AMDGPU.h:254
void emitKernel(const MachineFunction &MF, const SIProgramInfo &ProgramInfo) override
A parsed version of the target data layout string in and methods for querying it. ...
Definition: DataLayout.h:110
constexpr char Align[]
Key for Kernel::Arg::Metadata::mAlign.
Type
MessagePack types as defined in the standard, with the exception of Integer being divided into a sign...
Definition: MsgPackReader.h:48
raw_ostream & errs()
This returns a reference to a raw_ostream for standard error.
GCNRegPressure max(const GCNRegPressure &P1, const GCNRegPressure &P2)
This class represents an incoming formal argument to a Function.
Definition: Argument.h:29
LLVM_NODISCARD std::string str() const
str - Get the contents as an std::string.
Definition: StringRef.h:218
AMDGPU specific subclass of TargetSubtarget.
Address space for constant memory (VTX2).
Definition: AMDGPU.h:257
bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override
This class represents lattice values for constants.
Definition: AllocatorList.h:23
Calling convention for AMDGPU code object kernels.
Definition: CallingConv.h:200
A Module instance is used to store all the information related to an LLVM module. ...
Definition: Module.h:65
2: 32-bit floating point type
Definition: Type.h:58
amdgpu Simplify well known AMD library false FunctionCallee Value const Twine & Name
Document * getDocument() const
constexpr char PointeeAlign[]
Key for Kernel::Arg::Metadata::mPointeeAlign.
void toYAML(raw_ostream &OS)
Convert MsgPack Document to YAML text.
LLVM_NODISCARD bool startswith(StringRef Prefix) const
Check if this string starts with the given Prefix.
Definition: StringRef.h:256
bool hasFnAttribute(Attribute::AttrKind Kind) const
Return true if the function has the attribute.
Definition: Function.h:323
Metadata node.
Definition: Metadata.h:863
F(f)
const MDOperand & getOperand(unsigned I) const
Definition: Metadata.h:1068
uint64_t alignTo(uint64_t Value, uint64_t Align, uint64_t Skew=0)
Returns the next integer (mod 2**64) that is greater than or equal to Value and is a multiple of Alig...
Definition: MathExtras.h:684
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
Definition: DerivedTypes.h:534
1: 16-bit floating point type
Definition: Type.h:57
static IntegerType * getInt64Ty(LLVMContext &C)
Definition: Type.cpp:176
15: Pointers
Definition: Type.h:74
Address space for flat memory.
Definition: AMDGPU.h:253
StringSwitch & Case(StringLiteral S, T Value)
Definition: StringSwitch.h:67
Track resource usage for kernels / entry functions.
Definition: SIProgramInfo.h:21
A DocNode that is an array.
Type * getPointerElementType() const
Definition: Type.h:375
const DataLayout & getDataLayout() const
Get the data layout for the module&#39;s target platform.
Definition: Module.cpp:369
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
Definition: Twine.h:80
TypeID getTypeID() const
Return the type id for the type.
Definition: Type.h:137
constexpr char Printf[]
Key for HSA::Metadata::mPrintf.
Defines struct to track resource usage for kernels and entry functions.
MDNode * getMetadata(unsigned KindID) const
Get the current metadata attachments for the given kind, if any.
Definition: Metadata.cpp:1440
LLVM_NODISCARD R Default(T Value)
Definition: StringSwitch.h:181
bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override
LLVM_NODISCARD bool empty() const
empty - Check if the string is empty.
Definition: StringRef.h:126
bool fromYAML(StringRef S)
Read YAML text into the MsgPack document. Returns false on failure.
std::error_code fromString(std::string String, Metadata &HSAMetadata)
Converts String to HSAMetadata.
Key
PAL metadata keys.
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:244
static cl::opt< bool > VerifyHSAMetadata("amdgpu-verify-hsa-metadata", cl::desc("Verify AMDGPU HSA Metadata"))
DocNode getNode()
Create a nil node associated with this Document.
op_range operands() const
Definition: Metadata.h:1066
bool onlyReadsMemory() const
Return true if this argument has the readonly or readnone attribute.
Definition: Function.cpp:164
NamedMDNode * getNamedMetadata(const Twine &Name) const
Return the first NamedMDNode in the module with the specified name.
Definition: Module.cpp:250
constexpr char Attrs[]
Key for Kernel::Metadata::mAttrs.
AMDGPU HSA Metadata Streamer.
ArrayDocNode getArrayNode()
Create an empty Array node associated with this Document.
bool hasNoAliasAttr() const
Return true if this argument has the noalias attribute.
Definition: Function.cpp:133
11: Arbitrary bit width integers
Definition: Type.h:70
#define P(N)
A switch()-like statement whose cases are string literals.
Definition: StringSwitch.h:42
bool hasName() const
Definition: Value.h:250
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
The instances of the Type class are immutable: once they are created, they are never changed...
Definition: Type.h:45
static cl::opt< bool > DumpHSAMetadata("amdgpu-dump-hsa-metadata", cl::desc("Dump AMDGPU HSA Metadata"))
In-memory representation of kernel metadata.
ValueKind
Value kinds.
This file contains the declarations for the subclasses of Constant, which represent the different fla...
bool isPointerTy() const
True if this is an instance of PointerType.
Definition: Type.h:223
std::vector< uint32_t > mVersion
HSA metadata version. Required.
LLVM_NODISCARD size_t find(char C, size_t From=0) const
Search for the first character C in the string.
Definition: StringRef.h:285
Address space for private memory.
Definition: AMDGPU.h:259
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
constexpr char TypeName[]
Key for Kernel::Arg::Metadata::mTypeName.
amdgpu Simplify well known AMD library false FunctionCallee Value * Arg
Address space for region memory. (GDS)
Definition: AMDGPU.h:255
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function. ...
Definition: Function.cpp:196
static wasm::ValType getType(const TargetRegisterClass *RC)
static PointerType * getInt8PtrTy(LLVMContext &C, unsigned AS=0)
Definition: Type.cpp:219
MapDocNode getMapNode()
Create an empty Map node associated with this Document.
std::string & str()
Flushes the stream contents to the target string and returns the string&#39;s reference.
Definition: raw_ostream.h:498
Address space for local memory.
Definition: AMDGPU.h:258
SPIR_KERNEL - Calling convention for SPIR kernel functions.
Definition: CallingConv.h:136
unsigned getKernArgSegmentSize(const Function &F, unsigned &MaxAlign) const
AccessQualifier
Access qualifiers.
std::vector< std::string > mPrintf
Printf metadata. Optional.
unsigned getWavefrontSize() const
16: SIMD &#39;packed&#39; format, or other vector type
Definition: Type.h:75
CallingConv::ID getCallingConv() const
getCallingConv()/setCallingConv(CC) - These method get and set the calling convention of this functio...
Definition: Function.h:212
unsigned getParamAlignment() const
If this is a byval or inalloca argument, return its alignment.
Definition: Function.cpp:111
Module.h This file contains the declarations for the Module class.
LLVM_NODISCARD std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
Definition: StringRef.h:696
AddressSpace
Definition: NVPTXBaseInfo.h:21
unsigned getABITypeAlignment(Type *Ty) const
Returns the minimum ABI-required alignment for the specified type.
Definition: DataLayout.cpp:749
constexpr char CodeProps[]
Key for Kernel::Metadata::mCodeProps.
ValueType
Value types.
const Function & getFunction() const
Return the LLVM function that this machine code represents.
constexpr char LanguageVersion[]
Key for Kernel::Metadata::mLanguageVersion.
constexpr uint32_t VersionMinor
HSA metadata minor version.
unsigned getArgNo() const
Return the index of this formal argument in its containing function.
Definition: Argument.h:47
Simple in-memory representation of a document of msgpack objects with ability to find and create arra...
A DocNode that is a map.
constexpr char Kernels[]
Key for HSA::Metadata::mKernels.
void emitKernel(const MachineFunction &MF, const SIProgramInfo &ProgramInfo) override
uint64_t getTypeAllocSize(Type *Ty) const
Returns the offset in bytes between successive objects of the specified type, including alignment pad...
Definition: DataLayout.h:461
std::error_code toString(Metadata HSAMetadata, std::string &String)
Converts HSAMetadata to String.
This class keeps track of the SPI_SP_INPUT_ADDR config register, which tells the hardware which inter...
const Function * getParent() const
Definition: Argument.h:41
static const size_t npos
Definition: StringRef.h:50
unsigned getIntegerBitWidth() const
Definition: DerivedTypes.h:96
StringRef getValueAsString() const
Return the attribute&#39;s value as a string.
Definition: Attributes.cpp:194
constexpr char DebugProps[]
Key for Kernel::Metadata::mDebugProps.
AddressSpaceQualifier
Address space qualifiers.
StringRef getName() const
Return a constant reference to the value&#39;s name.
Definition: Value.cpp:214
int getIntegerAttribute(const Function &F, StringRef Name, int Default)
constexpr char AccQual[]
Key for Kernel::Arg::Metadata::mAccQual.
uint32_t Size
Definition: Profile.cpp:46
3: 64-bit floating point type
Definition: Type.h:59
virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict)=0
Emit HSA Metadata.
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
A raw_ostream that writes to an std::string.
Definition: raw_ostream.h:482
Module * getParent()
Get the module that this global value is contained inside of...
Definition: GlobalValue.h:565
Attribute getFnAttribute(Attribute::AttrKind Kind) const
Return the attribute for the given attribute kind.
Definition: Function.h:333
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:48
std::vector< Kernel::Metadata > mKernels
Kernels metadata. Required.
unsigned getNumOperands() const
Return number of MDNode operands.
Definition: Metadata.h:1074
const uint64_t Version
Definition: InstrProf.h:984
constexpr char Args[]
Key for Kernel::Metadata::mArgs.
constexpr uint32_t VersionMajor
HSA metadata major version.
iterator_range< arg_iterator > args()
Definition: Function.h:694