Add max threads checking for Metal #5588
Conversation
|
Xcode will do some checking like this: |
|
Please look at (and fix) the clang-tidy errors. |
Originally, this checking will be asserted by Metal API Validation in Xcode, otherwise the program will crash or output wrong results.
Thanks, fixed. |
| f.realize(output); | ||
| output.copy_to_host(); | ||
|
|
||
| for (int32_t i = 0; i < output.width(); i++) { |
There was a problem hiding this comment.
Is this loop necessary if we expect an error?
There was a problem hiding this comment.
It's not necessary indeed. Leave it just for code integrity...
|
Thanks for the PR! In principle, I'm not opposed to adding this check, though we don't do similar checks on any other backend. It might be better to only do the check in a debug target. I'd like to move towards a standard set of checks we do in non-debug and debug targets across all the GPU runtimes. Curious to hear what @zvookin thinks. |
| Var x("x"), y("y"); | ||
|
|
||
| f(x, y) = im(x, y) + 42; | ||
| f.gpu_blocks(y).gpu_threads(x, DeviceAPI::Metal); |
There was a problem hiding this comment.
What happens if the machine doesn't support Metal? I think this needs a check for whether the JIT target supports Metal, otherwise it should be skipped.
There was a problem hiding this comment.
Yeah, even if the machine supports metal, we should skip if the target doesn't support it, eg
if (!get_jit_target_from_environment().has_feature(Target::Metal)) {
printf("[SKIP] error/metal_threads_too_large ignored for targets without Metal enabled.\n");
_halide_user_assert(0);
}
There was a problem hiding this comment.
Thank you, since now it only checks in debug runtime, so I add a specific target, like "metal_vector_too_large" did.
|
@shoaibkamil Thanks, it's reasonable to only do this check in debug runtime, so I have moved it. For Cuda and OpenCL backends, they seem to do this checking in library. "cuLuanchKernel" and "clEnqueueNDRangeKernel" both have a return value to indicate this error: "cudaErrorLaunchOutOfResources" and "CL_INVALID_WORK_ITEM_SIZE". For OpenGL Compute backend, after "glDispatchCompute" is called, "GL_INVALID_VALUE" is generated to indicate this error, too. For D3D12Compute backend, since the threadgroups are encoded in the shader source, "D3DCompile" will failed when it exceeds the limit. So it seems the Metal backend is an exception. Not sure what behavior the Vulkan is. |
|
@shoaibkamil Do you also want to review before this lands? |
Originally, this checking will be asserted by Metal API Validation
in Xcode, otherwise the program will crash or output wrong results.