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

এক্সএলএ কাস্টম কল

এই নথিটি কীভাবে এক্সএলএর "কাস্টম কল" লিখতে এবং ব্যবহার করতে হয় তার বর্ণনা দেয়। কাস্টম কলগুলি আপনাকে একটি এক্সএলএ প্রোগ্রাম থেকে সি ++ বা সিইউডিএর মতো প্রোগ্রামিং ভাষায় লিখিত কোডটি চালু করতে দেয়।

সিপিইউতে কাস্টম-কল

আপনি একটি এইচএলও নির্দেশিকা তৈরি করতে পারেন যা এক্সএলএর ক্লায়েন্ট এপিআইয়ের মাধ্যমে একটি কাস্টম-কলকে উপস্থাপন করে। এটি লেখার মতো টেনসরফ্লো দিয়ে প্রকাশ করা হয়নি।

উদাহরণস্বরূপ, সিপিইউতে A[i] = B[i % 128] + C[i] গণনা করার জন্য নিম্নলিখিত কোডটি একটি কাস্টম-কল ব্যবহার করে। (অবশ্যই আপনি - এবং করা উচিত - নিয়মিত এইচএলও দিয়ে এটি করতে পারেন))

 #include "tensorflow/compiler/xla/client/xla_builder.h"
#include "tensorflow/compiler/xla/service/custom_call_target_registry.h"

void do_it() {
  xla::XlaBuilder b("do_it");
  xla::XlaOp param0 =
      xla::Parameter(0, xla::ShapeUtil::CreateShape(F32, {128}), "p0");
  xla::XlaOp param1 =
      xla::Parameter(1, xla::ShapeUtil::CreateShape(F32, {2048}), "p1");
  xla::XlaOp custom_call =
      xla::CustomCall(&b, "do_custom_call", /*operands=*/{param0, param1},
                      /*output_shape=*/ShapeUtil::CreateShape(F32, {2048}));
}

void do_custom_call(void* out, const void** in) {
  float* out_buf = reinterpret_cast<float*>(out);
  const float* in0 = reinterpret_cast<const float*>(in[0]);
  const float* in1 = reinterpret_cast<const float*>(in[1]);
  for (int i = 0; i < 2048; ++i) {
    out_buf[i] = in0[i % 128] + in1[i];
  }
}
XLA_REGISTER_CUSTOM_CALL_TARGET(do_custom_call, "Host");
 

লক্ষ্য করুন যে ফাংশনটি do_custom_call উপর চালিত বাফারগুলির মাত্রাগুলি জানতে হবে। এই উদাহরণে আমরা 128 এবং 2048 মাপকে হার্ডকোড করেছি you আপনি যদি এটি করতে না চান তবে আপনি কলগুলিতে প্যারামিটার হিসাবে মাত্রাটি পাস করতে পারেন।

জিপিইউতে কাস্টম-কল করুন

জিপিইউ কাস্টম কল ফ্রেমওয়ার্ক সিপিইউতে কিছুটা আলাদা। এখানে একটি সিইউডিএ উদাহরণ রয়েছে যা উপরের সিপিইউ কোড হিসাবে একই A[i] = B[i % 128] + C[i] গণনা করে।

 void do_it() { /* same implementation as above */ }

__global__ custom_call_kernel(const float* in0, const float* in1, float* out) {
  size_t idx = threadIdx.x * blockSize.x + gridIdx.x;
  out[idx] = in0[idx % 128] + in1[idx];
}

void do_custom_call(CUstream stream, void** buffers,
                    const char* opaque, size_t opaque_len) {
  const float* in0 = reinterpret_cast<const float*>(buffers[0]);
  const float* in1 = reinterpret_cast<const float*>(buffers[1]);
  float* out = reinterpret_cast<float*>(buffers[2]);

  const int64 block_dim = 64;
  const int64 grid_dim = 2048 / block_dim;
  custom_call_kernel<<<grid_dim, block_dim,
                       /*dynamic_shared_mem_bytes=*/0, stream>>>(in0, in1, out);
}
XLA_REGISTER_CUSTOM_CALL_TARGET(do_custom_call, "CUDA");
 

প্রথমে লক্ষ্য করুন যে জিপিইউ কাস্টম কল ফাংশনটি এখনও সিপিইউতে কার্যকর একটি ফাংশন । আমাদের do_custom_call CPU ফাংশনটি জিপিইউতে কাজ সারানোর জন্য দায়ী। এখানে এটি একটি CUDA কার্নেল চালু করে, তবে এটি কল কিউব্লাসের মতো অন্য কিছুও করতে পারে।

buffers পয়েন্টারের একটি অ্যারে যা হোস্টে থাকে এবং প্রতিটি উপাদান এটিতে ডিভাইস পয়েন্ট থাকে (যেমন জিপিইউ) মেমরি memory প্যারামিটারগুলি প্রথমে আসে, তারপরে আউটপুট মানটি অনুসরণ করে। এটি সিপিইউ কলিং কনভেনশন থেকে উল্লেখযোগ্যভাবে পৃথক, যার দুটি প্যারাম, ins এবং out । আমরা ডাইভার্জ করার মূল কারণটি হ'ল দক্ষতার সাথে টিউপল-আকারের ইনপুট / আউটপুটগুলি পরিচালনা করা; নীচের বিভাগটি দেখুন।

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

আপনাকে এটির চারপাশে কাজ করতে, আমরা opaque প্যারামিটার সরবরাহ করি। আপনি কাস্টম কলটি তৈরি করার সময় আপনি এটিগুলি একটি স্বেচ্ছাসেবী স্ট্রাইটে সেট করতে পারেন:

 std::string opaque = "...";
xla::CustomCall(&b, "do_custom_call", /*operands=*/{param0, param1},
                /*output_shape=*/ShapeUtil::CreateShape(F32, {2048}),
                opaque);
 

যেহেতু xla::Shape একটি প্রোটোকল বাফার উপস্থাপনা রয়েছে, আপনি এই ক্রমিক প্রোটোটিকে opaque অভ্যন্তরে সংরক্ষণ করতে পারেন এবং এটি আপনার জিপিইউ কাস্টম-কলের মধ্যে ডিসরিয়ালাইজ করতে পারেন। উল্লেখ্য তবে যে যদিও xla::ShapeProto ঘন ঘন পরিবর্তন করে না, এটি পরিবর্তন করে। অতীতে কীভাবে পরিবর্তন হয়েছে তা দেখতে গিট লগটি পরীক্ষা করুন।

কাস্টম-কলগুলিতে টিপলগুলি পাস করা

নিম্নলিখিত কাস্টম কল বিবেচনা করুন।

 using xla::ShapeUtil;
Shape p0_shape = ShapeUtil::MakeTuple({
    ShapeUtil::MakeShape(F32, {32}),
    ShapeUtil::MakeTuple({
        ShapeUtil::MakeShape(F32, {64}),
        ShapeUtil::MakeShape(F32, {128}),
    }),
    ShapeUtil::MakeShape(F32, {256}),
});
xla::XlaOp p0 = xla::Parameter(0, p0_shape, "p0");

Shape out_shape = ShapeUtil::MakeTuple({
  ShapeUtil::MakeShape(F32, {512}),
  ShapeUtil::MakeShape(F32, {1024}),
});
xla::CustomCall(&b, "do_custom_call", /*operands=*/{p0}, out_shape);
 

উভয় সিপিইউ এবং জিপিইউতে, একটি টুপল পয়েন্টারের অ্যারে হিসাবে মেমরিতে উপস্থাপন করা হয়। সি ++ - সিউডোকোডে, উপরোক্ত প্যারামিটার 0 নীচে উল্লেখ করা হয়েছে।

 // In-memory layout of parameter 0 from custom-call above.  True on both CPU
// and GPU.
float* subbuf0 = new float[32];
float* subbuf1 = new float[64];
float* subbuf2 = new float[128]
float* subbuf3 = new float[256];

void* subtuple = new void*[2];
(*subtuple)[0] = subbuf1;
(*subtuple)[1] = subbuf2;

void* p0 = new void*[3];
(*p0)[0] = subbuf0;
(*p0)[1] = subtuple;
(*p0)[2] = subbuf3;
 

যদিও সিপিইউ এবং জিপিইউতে টিউপসগুলির ইন-মেমরি উপস্থাপনা একই, তবুও তারা সিপিইউ এবং জিপিইউ কাস্টম-কল কলিং কনভেনশনগুলিতে আলাদাভাবে পরিচালিত হয়।

টেম্পল বাফার হিসাবে টুপল আউটপুট

কাস্টম-কলগুলিতে টুপল ইনপুটগুলি একটি সুবিধা, তবে সেগুলি কঠোরভাবে প্রয়োজন হয় না। আমরা যদি কাস্টম কলগুলিতে টুপল ইনপুটগুলিকে সমর্থন না করি, আপনি কাস্টম কলটিতে যাওয়ার আগে টিপলগুলি get-tuple-উপাদান ব্যবহার করে সর্বদা আনপ্যাক করতে পারেন।

অন্যদিকে, টুপল আউটপুটগুলি আপনাকে এমন কাজ করতে দেয় যা আপনি অন্যথায় করতে পারেন নি।

টুপল আউটপুট থাকার সুস্পষ্ট কারণ হ'ল এটিই কীভাবে কাস্টম কল (বা অন্য কোনও এক্সএলএল অপ) একাধিক স্বতন্ত্র অ্যারে প্রদান করে।

তবে স্পষ্টতই, একটি কচি আউটপুটও আপনার কাস্টম কল টেম্প মেমোরি দেওয়ার একটি উপায়। হ্যাঁ, একটি আউটপুট একটি অস্থায়ী বাফার উপস্থাপন করতে পারে। বিবেচনা করুন, একটি আউটপুট বাফারের সেই সম্পত্তি রয়েছে যা ওপেন এটিতে লিখতে পারে এবং এটি লিখিত হওয়ার পরে এটি এটি থেকে পড়তে পারে। টেম্প বাফার থেকে আপনি যা চান তা ঠিক এটি।

উপরের উদাহরণে, ধরুন আমরা একটি টেম্প বাফার হিসাবে F32[1024] ব্যবহার করতে চেয়েছিলাম। তারপরে আমরা ঠিক উপরের মতো এইচএলও লিখতে চাই এবং আমরা কাস্টম কলের আউটপুটটির টিউপল সূচক 1 কখনই পড়তে চাই না।

সিপিইউ কাস্টম-কলগুলিতে টিপলস

সিপিইউ কোডে, আমাদের একটি ক্রিয়াকলাপ do_custom_call(const void** ins, void* out)ins কেবল একটি উপাদান সহ একটি অ্যারে, যা param0 । প্যারাম0 এর param0 সেই পয়েন্টারটিকে output_tuple অ্যাক্সেসযোগ্য এবং আউটপুট_টুপল এর output_tuple out করে অ্যাক্সেসযোগ্য।

জিপিইউ কাস্টম-কলগুলিতে টিপলস

জিপিইউ কোডে, আমাদের একটি ক্রিয়া do_custom_call(..., void** buffers, ...) । এই ক্ষেত্রে buffers ছয়টি ডিভাইস পয়েন্টারের হোস্ট অ্যারে, ইনপুট / আউটপুটটিতে প্রতিটি লিফ বাফারের জন্য একটি। ফ্ল্যাট তালিকাটি তৈরি করতে, আমরা পরামিতিগুলি এবং আউটপুট ধরে পুনরাবৃত্তি করি এবং প্রত্যেকটির জন্য আমরা এর আকারটির প্রাক-অর্ডার ট্রভারসাল করি। মূর্তভাবে:

 // Layout of `buffers` parameter to GPU custom call function for custom-call
// above.
buffers[0] == subbuf0
buffers[1] == subbuf1
buffers[2] == subbuf2
buffers[3] == subbuf3
buffers[4] == output_subbuf0
buffers[5] == output_subbuf1