MCPcopy Create free account
hub / github.com/arrayfire/arrayfire / passesJitHeuristics

Function passesJitHeuristics

src/backend/oneapi/Array.cpp:320–414  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

318/// limitation on the platform. For NVIDIA this is 4096 bytes. The
319template<typename T>
320kJITHeuristics passesJitHeuristics(span<Node *> root_nodes) {
321 if (!evalFlag()) { return kJITHeuristics::Pass; }
322 static auto getLogger = [&] { return common::loggerFactory("jit"); };
323 for (const Node *n : root_nodes) {
324 if (n->getHeight() > static_cast<int>(getMaxJitSize())) {
325 AF_TRACE(
326 "JIT tree evaluated because of tree height exceeds limit: {} > "
327 "{}",
328 n->getHeight(), getMaxJitSize());
329 return kJITHeuristics::TreeHeight;
330 }
331 }
332
333 // TODO(umar): add memory based checks for JIT kernel generation
334 bool isBufferLimit =
335 false; // getMemoryPressure() >= getMemoryPressureThreshold();
336 // auto platform = getActivePlatform();
337
338 // The Apple platform can have the nvidia card or the AMD card
339 // bool isIntel = platform == AFCL_PLATFORM_INTEL;
340
341 /// Intels param_size limit is much smaller than the other platforms
342 /// so we need to start checking earlier with smaller trees
343 int heightCheckLimit = 3;
344
345 // A lightweight check based on the height of the node. This is
346 // an inexpensive operation and does not traverse the JIT tree.
347 bool atHeightLimit =
348 std::any_of(std::begin(root_nodes), std::end(root_nodes),
349 [heightCheckLimit](Node *n) {
350 return (n->getHeight() + 1 >= heightCheckLimit);
351 });
352
353 if (atHeightLimit || isBufferLimit) {
354 // This is the base parameter size if the kernel had no
355 // arguments
356 size_t base_param_size =
357 (sizeof(T *) + sizeof(Param<T>)) * root_nodes.size() +
358 (3 * sizeof(uint));
359
360 const sycl::device &device = getDevice();
361 size_t max_param_size =
362 device.get_info<sycl::info::device::max_parameter_size>();
363 // typical values:
364 // NVIDIA = 4096
365 // AMD = 3520 (AMD A10 iGPU = 1024)
366 // Intel iGPU = 1024
367 max_param_size -= base_param_size;
368
369 struct tree_info {
370 size_t total_buffer_size;
371 size_t num_buffers;
372 size_t param_scalar_size;
373 };
374
375 tree_info info{0, 0, 0};
376 for (Node *n : root_nodes) {
377 NodeIterator<> it(n);

Callers

nothing calls this directly

Calls 9

loggerFactoryFunction · 0.85
beginFunction · 0.85
endFunction · 0.85
getDeviceFunction · 0.85
getHeightMethod · 0.80
isBufferMethod · 0.80
getBytesMethod · 0.45
getParamBytesMethod · 0.45

Tested by

no test coverage detected