This page was translated by the Cloud Translation API.
Switch to English

XLA জন্য MLIR CodeGen

XLA করে পরিচালিত HloInstruction এবং এই প্রতিনিধিত্ব অনেক অপ্টিমাইজেশন সঞ্চালিত, লক্ষ্যবস্তু ডিভাইসের মধ্যে এই অনেক ভাগ করুন। কিছু পয়েন্ট হিসাবে একটি রৈখিক সময়তালিকা নির্ণয় করা হয় এবং মেমরি বাফার স্ট্যাটিক্যালি প্রতিটি মান নির্ধারিত হয়। ডিভাইস নির্দিষ্ট codegen এই ক্রম ঢোঁড়ন এবং "এমিটার্স" কলিং একটি উপস্থাপনা ডিভাইসের জন্য উপযুক্ত (উদাহরণস্বরূপ CPU তে XLA গণনার প্রতি একটি একক LLVM ফাংশন, বা "thunks" একটি ক্রম জিপিইউ অপারেশন encapsulating জেনারেট করতে এবং সম্ভবত উত্পন্ন PTX যখন দ্বারা পরিচালনা লক্ষ্য করে জিপিইউ)।

একটি উপস্থাপনকারী ধাপ হিসেবে, প্রক্রিয়া ডান আটকাচ্ছে পর XLA বাফার-নিয়োগ ফেজ সমাপ্ত প্রক্রিয়ায় বর্তমানে হন এবং পরিবর্তে নির্গত একটি MLIR মডিউল lhlo উপভাষা। থেকে সেখানে আমরা (প্রধানত Linalg, অ্যাফিন, ও GPU উপভাষা) MLIR উপাদান ব্যবহার করে codegen সঞ্চালন ডিভাইস উপর নির্ভর করে।

নীচে রেকর্ডের পরিকল্পনা বৃদ্ধিলাভ XLA / জিপিইউ মাইগ্রেট করতে ব্যবহার করা lhlo codegen ইনপুট হিসাবে।

কাজ

নিমন্ত্রণকর্তা যন্ত্র
ছক পূরণ করা HloInstruction * (টাস্ক 1) HloInstruction * (টাস্ক 1)
আউটপুট ফরমেট xla :: Thunk (টাস্ক 2) LLVM আইআর (টাস্ক 3)
  • কার্য 1 LHLO করার HloInstruction * থেকে উভয় হোস্ট এবং ডিভাইস ইনপুট বিন্যাস পরিবর্তন।
  • কার্য 2 "হোস্টের জন্য কিছু অবতরণ প্যাড" (নীচে দেখুন) থেকে thunks থেকে হোস্ট আউটপুট বিন্যাস পরিবর্তন।
  • MLIR কিছু ফর্ম থেকে LLVM আইআর থেকে টাস্ক 3 মাইগ্রেট ডিভাইস আউটপুট। এটা এই প্রকল্পে ঐচ্ছিক, এবং বিস্তারিত জানার জন্য অধ্যায় "মাইগ্রেট ডিভাইস LLVM আইআর" দেখুন।

এই প্রকল্পটি LHLO-এমিটার্স সঙ্গে এন্ড-টু-এন্ড runnable মডেলের থাকার অগ্রাধিকার দেয় যতটা সম্ভব সক্ষম করা হয়েছে। এর অর্থ হলো অগ্রাধিকারের ভিত্তিতে উদ্দেশ্য নিম্নলিখিত ক্রমে তালিকা:

  • বিদ্যমান Thunks এবং এমিটার্স অপরিবর্তিত সঙ্গে XLA / LHLO এমিটার্স সঙ্গে জিপিইউ runnable করুন।
  • ক্ষেত্রে দ্বারা LHLO মধ্যে HloInstruction *, ক্ষেত্রে রেফারেন্স বর্জন করুন:
    • একটি MLIR ভিত্তিক বিকিরণকারী করার জন্য একটি উত্তরাধিকার বিকিরণকারী স্যুইচ করুন (যেমন Linalg), অথবা
    • মেকানিক্যালি MLIR উপস্থাপনা নিতে (জিপিইউ ডায়ালেক্ট সঙ্গে স্ট্যান্ডার্ড স্থানান্তর) বিদ্যমান বিকিরণকারী অনুবাদ করুন।

মাইগ্রেট Thunks (টাস্ক 2)

xla :: GPU :: Thunk একটি ডাটা স্ট্রাকচার যে:

  • হোস্ট থেকে বলা যেতে পারে (xla :: GPU :: Thunk :: ExecuteOnStream ())।
  • তার উপশ্রেণী বিভিন্ন তথ্য বহন করে।
  • BufferAllocation :: স্লাইস এবং StreamExecutor সাথে মিথস্ক্রিয়া।
  • উত্ক্ষেপণের কার্নেলের
  • সব রানটাইম লাইব্রেরি মধ্যে কল।

যে খরচ রয়েছে:

  • অপ-নির্দিষ্ট কনফিগারেশন ডেটা প্রতিনিধিত্ব (যেমন সংবর্তন configs)।
  • অপ আকৃতি এবং প্রতীক আকার মাইগ্রেট।
  • thunks (একই শর্তে, ইত্যাদি) একটি গাছ প্রতিনিধিত্ব।

মাইগ্রেশন কাজ LHLO / বিকিরণকারী মাইগ্রেশন থেকে স্বাধীন। সীমিত সম্পদের অধীনে, এটা LHLO পিছনে অগ্রাধিকারের এর / মাইগ্রেশন বিকিরণকারী।

আমরা কিভাবে LHLO থেকে হোস্ট প্রান্তের অংশ নীচু কাছে অনেক বিকল্প আছে:

  • TFRT
    • ব্যবহারের জন্য (প্রো) মহান CUDA এবং নিতম্ব চাদরে।
    • (পূর্বে-) সহজ, গ্রন্থাগার কল (cuDNN, cuBLAS, cuFFT, ইত্যাদি) বাস্তবায়ন যেমন TFRT অপস সি ++ কোড দ্বারা ব্যাখ্যা করা হয়।
    • (কন) হোস্ট পাশ বিকাশ চলছে এবং পরীক্ষিত নয়।
  • Jitted CPU- র কোড
    • (পূর্বে-) মহান কম ক্ষমতা। কয়েক loops এবং অবস্থার তৈরি করুন এবং এটি সম্পন্ন হয়েছে।
    • (কন) GPUDialect এখনো মডেল চেইন / স্ট্রিম / asynchronicity / ডিভাইস বরাদ্দ আছে।
    • (কন) CUDA / হিপ রানটাইম সমর্থন ন্যূনতম (টুলকিট পথ, সংস্করণ, গতিশীল লোডিং, ইত্যাদি) হয়।
  • বর্তমান (ব্যাখ্যা) XLA রানটাইম

ডিসিশন: TFRT দত্তক গ্রহণ করা, কিন্তু TFRT মধ্যে jitting CPU- র কোড সমর্থন করি।

মাইগ্রেট ডিভাইস LLVM আইআর (টাস্ক 3)

একটি আধিভৌতিক বিকিরণকারী উপাদান দ্বারা এটি উপাদান পূরণ করে লক্ষ্য অপ জেনারেট করে। প্রতিটি আউটপুট উপাদান operands থেকে উপাদানগুলি একটি সেট উপর নির্ভর করে। সমস্ত উপাদান গতিশীল সূচকের সঙ্গে বাফার মিশ্রন দ্বারা বর্ণনা করা হয়। এটা প্রায় সমস্ত "গণিত" অপস বর্ণনা করতে যথেষ্ট, কিন্তু কর্মক্ষমতা কারণে কেবল "গণিত" অপস বৃহৎ উপসেট মধ্যে (CPU | GPU) সরাসরি প্রয়োগ করা হয় ElementalIrEmitter।

ElementalIrEmitter যে অদ্বিতীয়:

  • কোডের একটি বড় অংশ XLA / জিপিইউ এবং CPU- র মধ্যে ভাগ করা আছে।
  • এটা তোলে মডেল দেখা অপস, সমস্ত উপাদান-অনুযায়ী অপস সহ একটি বৃহৎ অংশ প্রতিনিধিত্ব করে।
  • সর্বাধিক fusions একমাত্র ElementalIrEmitter উপর নির্ভর করে।
  • যেমন অপ উপাদান এবং প্রতীক উপাদানের মধ্যে একটি ডাটা নির্ভরতা DAG বর্ণনা এটা গঠনের দিক সহজ।
  • এটা বেশিরভাগ পোর্টেবল এবং উচ্চ পর্যায়ের (যেমন জিপিইউ kReduce ও GPU kCopy অসদৃশ)।
  • ডায়নামিক আকৃতি সমর্থন সহজ অন্তত উপাদান-অনুযায়ী অপস জন্য।

এখন, সব অপস জন্য, elementally-নির্গত বা না, প্রতিটি XLA অপ শেষে রাজ্যের বিভিন্ন স্বাদে আছে:

  1. LLVM আইআর হিসেবে ডিভাইস কোড থাকবে।
  2. পুরাতন বিকিরণকারী Refactor LHLO মত হতে -> MLIR LLVM ডায়ালেক্ট:
    • (খরচ) থ্রো-দূরে কাজ করা হবে যদি আমরা শেষ পর্যন্ত স্ট্যান্ডার্ড স্থানান্তর করতে চান।
    • (বেনিফিট) এটা সহজ এবং যান্ত্রিক হয়। অল্প সময়ের মধ্যে সম্পন্ন করা যাবে না।
    • (বেনিফিট) এটা (1) তুলনায় আরো ফলপ্রসূ হবে না।
  3. পুরাতন এমিটার্স Refactor LHLO মত হতে -> MLIR জিপিইউ + + স্ট্যান্ডার্ড + + Loops:
    • (খরচ) উদ্ধরণ বিদ্যমান স্ট্যান্ডার্ড এমিটার্স কিছু চ্যালেঞ্জ প্রবর্তন করে। পয়েন্টার এবং GEPs MemRefs এবং SubViews রূপান্তরিত করা প্রয়োজন। amdgpu সম্পূর্ণতার সুনিশ্চিত অন্য এক।
    • (খরচ) XLA / জিপিইউ প্রচন্ডভাবে LLVM মেটাডাটা ভরসা করে রয়েছে:
      • range ব্লক / থ্রেড সূচকের জন্য।
      • align , dereferenceable , invariant.load , alias.scope , noalias লোড / দোকানে।
      • llvm.loop.unroll.disable , llvm.loop.unroll.full , llvm.loop.vectorize.enable অনুক্রমিক loops জন্য।
    • (বেনিফিট) দীর্ঘমেয়াদী হতে পারে। আরো পোর্টেবল।
  4. > Linalg, এবং নতুন Linalg এমিটার্স লিখুন - Refactor পুরাতন এমিটার্স LHLO হতে
    • (খরচ) এই ক্ষেত্রে দ্বারা ক্ষেত্রে দেখা যায়। পূর্ববর্তী অপশন তুলনায় একটি নতুন বাস্তবায়ন যে XLA এর পারফরম্যান্সের সাথে মিলে যায় বেঞ্চমার্ক মধ্য দিয়ে যেতে প্রয়োজন <-> অপ্টিমাইজ কর্মপ্রবাহ, যা কিছু অপস জন্য একটি উল্লেখযোগ্য খরচ হতে পারে।
    • (বেনিফিট) ইউনিফাইড স্ট্যাকের; সম্প্রদায় সমর্থন; বহনযোগ্যতা; আরো অপ্টিমাইজেশান সম্ভাবনা।

উপসংহার:

  • (2) জন্য যেতে করবেন না। (1) বা (3) শুধু (2) চেয়ে ভাল। (2) খরচ চেয়ে বেশি (1), যেহেতু এটি যান্ত্রিক refactoring অনেকটা প্রয়োজন। সঙ্গে (1) আমরা এখনও XLA সক্রিয় MLIR এমিটার্স নিতে লক্ষ্য অর্জন করতে পারেন। > LLVM আইআর - -> Run উত্তরাধিকার ডিভাইস এমিটার্স এই LHLO করছেন হয়।
  • ElementalIrEmitter অপস (4) জন্য যেতে, কিন্তু বৃদ্ধিলাভ না। কারণ সব elementally-নির্গত অপস একই গ্রাফ মধ্যে সংযুক্ত আছেন সেখানে, অপ দ্বারা এটি অপ করার কোন উপায় নেই। (কার্নেল জেনারেটর, Linalg xla / সেবা / mlir_gpu) এই কাজ বিভিন্ন চালু বাহিনীর একটি একীকরণ পয়েন্ট হিসাবে পরিবেশন করতে পারেন।
  • অন্যান্য সকল অপস জন্য (1) যান। একটি প্রসারিত লক্ষ্য হিসাবে, তারা (3) বা (4) চলে আসেন হতে পারে।

অগ্রাধিকার করণ

যদিও এই তিনটি কর্ম উপরে উল্লিখিত parallelizable হয়, সীমিত সম্পদের অধীনে তারা ধারাবাহিকভাবে করতে হবে। অগ্রাধিকার প্রতিটি কার্যের সমাপ্তির জন্য দৃশ্যমান ফলাফল উপর গুরুত্ত্ব দেয়।

অগ্রাধিকার হল: Task1 (উত্তরাধিকার এমিটার্স জন্য LHLO)> টাস্ক 2 (Thunks)> টাস্ক 3 (MLIR এমিটার্স)।

কার্য 1 শেষে XLA ব্যবহারকারীদের একটি LHLO উৎপন্ন (যেমন কার্নেল জেনারেটর) এবং তাদের নির্বাহ করতে পারেন। সংকলন বিন্যাস serializable MLIR হবে না।

কার্য 2 শেষে LHLO সঠিক, serializable MLIR করার কমে যায়। এই অফলাইন সংকলন দেয়।

কার্য 3 শেষে সব XLA এমিটার্স তার বাস্তবায়ন মধ্যে MLIR ভিত্তিক হয়।

বিস্তারিত নকশা

ধাপ 1: (টাস্ক 1) সমাপ্তি LHLO এবং উত্তরাধিকার এমিটার্স করুন LHLO নিন

এই ধাপে MLIR অপস সাথে ইন্টারঅ্যাক্ট সমস্ত বিদ্যমান XLA / জিপিইউ এমিটার্স করে তোলে। এই ধাপে বিশুদ্ধ refactoring এবং NFC হয়।

এই ধাপে বেশিরভাগ যান্ত্রিক, কিন্তু এটা একটি unnested HloComputation এবং LHLO মধ্যে নিম্নলিখিত গোলযোগ ঠাহর দরকারী নয়:

  • প্রতিটি HloInstruction তার operands থেকে সরাসরি প্রবেশাধিকার (ক ডেটা-প্রবাহ DAG) আছে। বিপরীতভাবে, প্রতিটি LHLO অপ শুধুমাত্র তার প্রতীক বাফার অ্যাক্সেস (অপস এবং বাফার মধ্যে একটি দ্বিপাক্ষিক) আছে। LHLO অপস তাদের প্রতীক অপস অ্যাক্সেস করতে মাধ্যমে ব্যবহার-Def চেইন যেতে হবে।
  • Unnested উত্তরাধিকার এমিটার্স প্রায়োগিক প্রায় তাদের operands অ্যাক্সেস না। একমাত্র ব্যতিক্রম kReduce হয়।
  • Unnested উত্তরাধিকার এমিটার্স, শুধুমাত্র পেয়ে টুকরা জন্য অ্যাক্সেসের BufferAssignment dataflow_analysis () বা alias_analysis মত-aux-ডাটা স্ট্রাকচার ব্যবহার করার জন্য নয় ()। llvm_ir নিজস্ব alias_analysis () ফালি তথ্যের উপর ভিত্তি করে তৈরী করে।

উপসংহার LHLO প্রধান অসুবিধা ছাড়াই ডান-মাপসই হবে।

পদক্ষেপ 2: (ঐচ্ছিক) প্রোফাইলিং সাপোর্ট

এই ধাপে শুধুমাত্র প্রয়োজন হয় যদি আমরা XLA Thunk যুক্তিবিজ্ঞান কিছু পরিত্যাগ করতে (পরবর্তী পদক্ষেপ দেখুন) শুরু।

আসলে কোনো MLIR ভিত্তিক এমিটার্স চালু করার আগে, আমরা MLIR ভিত্তিক এমিটার্স জন্য প্রোফাইলিং প্রয়োজন।

বর্তমানে XLA StreamExecutor এর টাইমার মধ্যে কল করে তার নিজস্ব প্রোফাইলিং সম্পাদন করে। ফণা অধীন টাইমার দুটি ঘটনার আগে ও একটি কার্নেল প্রবর্তন পরে সন্নিবেশ এবং পরিমাপ করে এই দুটি ইভেন্টের মধ্যে সিঙ্ক সময়।

সেখানে প্রায় তিন পন্থা MLIR মধ্যে প্রোফাইলিং সমর্থন করার জন্য আছেন:

  • একটি প্রোফাইলার এন্ড-টু-এন্ড চালান
  • LHLO প্রতিটি অপ জন্য একটি প্রোফাইল অপ যোগ করুন, একটি ইনজেকশনের প্রোফাইলার ব্যবহার করে।

"এন্ড-টু-এন্ড" অভিগমন MLIR স্বচ্ছ, কিন্তু একই সমস্যা যে XLA তোলে প্রথম স্থানে এটি ব্যবহার করবেন ভুগছেন: একটি প্রোফাইলার দ্বারা সংগৃহীত গ্রন্থাগার কল (nvprof / ...) সহজে HLO সঙ্গে সম্পর্কযুক্ত করতে পারবে না অপস। উদাহরণ হিসেবে বলা যায়, প্রতিটি HLO জন্য cuDNN লঞ্চ একাধিক কার্নেল, এবং এটি কঠিন বলতে যা যা HLO মিলা কার্নেলের।

"ইনজেকশনের প্রোফাইলার 'পদ্ধতির প্রয়োজন:

  • LHLO একটি প্যারামিটার হিসেবে প্রোফাইলার নিতে।
  • সামনে এবং প্রতিটি অপ পর profile.start / profile.end ঢোকাতে।
  • একটি যে পরিধেয় প্রোফাইল থেকে পাস। একটি সি ++ বাস্তবায়নে {শুরু, শেষ}।

সঠিক প্রোফাইলিং সহজে MLIR-জেনারেট অপস জন্য সম্পন্ন করা যাবে না, যেহেতু:

  • MLIR একটি টাইমার নেই, না এটা TFRT / StreamExecutor উপর নির্ভর করে।
  • MLIR সহজে জটিল পরামিতি সঙ্গে সি ফাংশন মধ্যে কল করে না।

ধাপ 3: (টাস্ক 2) মাইগ্রেট Thunks

একটি নোট হিসাবে, সেখানে thunks প্রায় তিন প্রকারঃ আছেন:

  • KernelThunk, যা একটি কার্নেল আরম্ভ করা হয়।
  • নিয়ন্ত্রণ প্রবাহ thunks, যা (, শর্তাধীন সময়, জন্য, ক্রম) হোস্ট নিয়ন্ত্রণ প্রবাহ যুক্তিবিজ্ঞান এবং লঞ্চ শরীর কার্নেলের।
  • লাইব্রেরী thunks: cuDNN, cuBLAS, cuFFT, NCCL, ইত্যাদি

পরিকল্পনা:

  • Thunks (ডি) serializable করুন।
  • সাহায্য একটি রাষ্ট্র যেখানে এটি এই শব্দার্থবিদ্যা সমর্থন করতে TFRT উন্নত।
  • রাষ্ট্র উন্নত হিসাবে, পৃথক thunks বৃদ্ধিলাভ মাইগ্রেট।

এই কর্ম আইটেমগুলি আংশিকভাবে আদেশ হয়। প্রকৃত মৃত্যুদন্ড অর্ডার / প্রকৌশল উপমা যেমন যায় মূল্যায়ন করা হয়।

ধাপ 4: (টাস্ক 3) মাইগ্রেট করা ElementalIrEmitter

একবার প্রোফাইলিং প্রস্তুত, আমরা সম্পূর্ণ সুর MLIR সমস্ত ElementalIrEmitter ভিত্তিক এমিটার্স পারবেন না। তারপর আমরা ডিফল্টরূপে তাদের চালু করেন, তখন অভিমানী যে এই MLIR ভিত্তিক এমিটার্স সব একটি একক প্রবাহ ব্যবহার করুন।

লক্ষ করুন যে XLA / CPU- এর ElementalIrEmitter মাইগ্রেট পাশাপাশি, যেহেতু তারা কোডের একটি বড় অংশ ভাগ উপকারী নয়।

সঙ্গে সব মাপকাঠিতে এবং কর্মক্ষমতা শিকার করা সম্পন্ন (করণীয়: কর্মক্ষমতা সমতা সংজ্ঞায়িত), আমরা নতুন MLIR ভিত্তিক মৌল বিকিরণকারী চালু করেন, তখন ও উত্তরাধিকার ElementalIrEmitter মুছে দিন।

এই ধাপে পরে মাইগ্রেশন জন্য সহজ লয় ট্রানজিশন (নেস্টেড অপস) প্রদান করে।

পদক্ষেপ 5: মাল্টি স্ট্রিম সাপোর্ট অথবা ড্রপ

আমরা মুছে ফেলতে পারবেন না এমিটার্স কিছু যতক্ষণ না আমরা MLIR মধ্যে এটিকে সমর্থন করে, অথবা আমরা বৈশিষ্ট্য ড্রপ। এটা তোলে MLIR কাজ অপেক্ষাকৃত বড় পরিমাণ এবং XLA জন্য লাভ অল্প পরিমাণ আছে। আমরা বহু-স্ট্রীম XLA / জিপিইউ ব্যবহারকারীদের বর্তমান ব্যবহারকারীদের তদন্ত, এবং যদি যুক্তিসঙ্গত এই বৈশিষ্ট্যটি মুছে ফেলতে চেষ্টা করা উচিত।

ধাপ 6: (টাস্ক 3) মাইগ্রেট করা ডিভাইস অপস

এই ধাপে সব unnested অপস মাইগ্রেট, তাহলে আমরা সব unnested এমিটার্স মুছে দিতে পারেন।

kCopy এবং kReduce জন্য একটি লেখা / refactor এই কল। kReduce ইতিমধ্যে প্রচুর জন্য কাজ তাই কাজ প্রকৃত পরিমাণ যে চাহিদা কাজ করতে হবে অবশেষ দেখা হবে হয়।