| 318 | /// limitation on the platform. For NVIDIA this is 4096 bytes. The |
| 319 | template<typename T> |
| 320 | kJITHeuristics 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); |
nothing calls this directly
no test coverage detected