iop_profile: never alias source and target of a GPU colorspace conversion - #21848
iop_profile: never alias source and target of a GPU colorspace conversion#21848piratenpanda wants to merge 3 commits into
Conversation
piratenpanda
changed the title
jenshannoschwalm
added
OpenCL
|
I am very well aware of the 1.2 OpenCL specs and the mis-use of providing the same cl_mem as input and output (alias) of a kernel and it's not the first time discussing this :-) But - we have been using such code at a number of places in dt and i could never (yet?) track down an issue related to this use-against-the-standard on any driver if we strictly read&write just a single location (i think that's what above "Being pointwise") (Maybe that's related to the hardware cacheline ...) Just noting - it's only finalscale that outputs NaN in the log - do you understand why that's not detected earlier? If we/you think this against-the-specs code is something that has to go, ok. (We might change code is using OpenCL 3.0) About this PR, if we/you want to go according to specs, we should not process&backcopy-from scratch but ensure we process to a new cl_mem image and use that later if possible to avoid a costly image-copy. (From memory that would be pretty straighforward ...) Also remember, that we might have to check if available cl_mem might be stressed (just alloc and test vs non-NULL is not enough) |
|
I have no hard feelings about this and as said, I have no idea about opencl. I just see lots of blocks in the image after zooming in with HQ enabled before this patch and none after this patch. Here's a reproducer for me: https://discuss.pixls.us/t/spektrafilm-darktable-module-implementation-discussion/58744/486?u=piratenpanda Enable HQ processing, zoom in, see this: |
|
Maybe it's also due to Arch's: |
|
So, why not do it without the internal scratch for performance? Or would you want me to do that? |
|
Aah, i have seen other reports about that lately. |
|
Seems to be a rusticl issue after all. @karolherbst is this maybe already fixed in mesa 26.2? Tests created by Claude: Reproducer
It runs one pointwise kernel over identical input for every combination of image size, sampler addressing mode, and aliased vs separate target, counting values that differ from the expected result. __constant sampler_t samp = CLK_NORMALIZED_COORDS_FALSE
| CLK_ADDRESS_NONE
| CLK_FILTER_NEAREST;
__kernel void convert(__read_only image2d_t in, __write_only image2d_t out,
const int width, const int height)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
if(x >= width || y >= height) return;
float4 p = read_imagef(in, samp, (int2)(x, y));
write_imagef(out, (int2)(x, y), p * 2.0f + 1.0f);
}Images are ResultsDamage lands in the trailing rows, and only when the target aliases the source. Height scan at width 1024: Wrong values include Observations
That the damage is a pure function of image dimensions, unchanged by workgroup shape or kernel, seems to point at resource handling for a partial tile at the bottom edge rather than at anything in shader execution. Darktable's sidedarktable's OpenCL pipeline converts image buffers between colour spaces in place, passing one darktable's buffer in the observed failure is 3979x5817, one of the sizes the reproducer above finds corrupt. This still does not make sure that the two are the same bug. The reproducer damages 2360 pixels confined to the final row, none of them isolated. darktable's corrupted buffer, at the same dimensions and through the same in-place aliasing, has 14351 damaged pixels spread across 2118 rows, 98% of them isolated single pixels. Same trigger, different signature, and the difference is unexplained. The untested differences are that darktable's conversion kernel binds a second |
|
@jenshannoschwalm It also seems that me changing the pipe position of the Monochrome module triggers this in the first place. If I move it behind the tonemappers as I did to get rid of spektrafilms color film noise this causes the nans. Maybe it's just "user error" in this case? |
|
Dont know, there is a pending PR btw for better NaN checks btw. |
|
With your PR: |
Yeah.. Using the same image as Quote from the spec for clSetKernelArg
If you want to read from and write to the same image object you'll have to use |
|
Thanks @karolherbst for stepping in here. I confess that i don't understand the aliased case situation but i conclude from above that current dt code "just works" in many cases but fails on certain "dimensions" / drivers. I conclude that we will have to go the hard way -- changing all dt code where we don't follow the specs ... |
Yeah.. the thing is, that drivers might heavily cache or prefetch I think practically it should be fine if all reads are done prior to the writes to the same locations, but I don't think it's a good idea to rely on implementation defined behavior either way. And that also gets more complicated with filtering. There might be a bug still, but if it only shows with images used for |
|
@piratenpanda THANKS A LOT for your sweep test! @TurboGit i checked the proposed commit again and it's callers. Yes - it's safe and dt assumptions about used cl_mem are not hurt. So i would say we should use it for master and 5.6.1. and i will later take care for performance and other dt code parts where we do badly. |
I'll let Claude run an overview about this later and then maybe you can take this as a comparison/starting point when you are back |
That's exactly - at least how i understand it - what i/we assumed. Maybe this is - the only-sometime-failing-tests might explain this - the device hardware "cacheline" width? |
|
As far as I understood the logs the NaNs appear when the Cache is hit, yes |
jenshannoschwalm
added
the
priority: high
|
will rework this |
8b402e1 to
df265e4
Compare
|
redone according to comments, split into three commits for the respective modules |
|
Checked, that's all just perfect :-) The atrous tiling requirement is correct, the pixelpipe code is also perfect with minimally perf cost, lut3d also. Just tested integration suite - to me some subtle diffs but nothing to be worried about. The iop_profile commit would need some mods for 5.6.1 |
I think the problem of relying on this to be safe is also making assumptions of GPU internal behavior, so who knows. Could be cachelines, could be some other GPU internal implementation detail. |
4 participants
While playing around with the strange history stack I created with the nested style bug I found another case of NaNs in the image when running the HQ pipe with rusticl:
opencl_log.txt
Tagging @jenshannoschwalm. This fixes the issue for me but I am no opencl expert. According to Claude's output OpenCL spec 1.2 §6.12.14 says a kernel cannot both read from and write to the same image object which makes this a darktable bug rather than a rusticl one.
From Claude:
dt_ioppr_transform_image_colorspace_cl() and dt_ioppr_transform_image_colorspace_rgb_cl() passed the caller's buffer
as both the read_only source and the write_only target of the conversion kernels whenever the caller handed in one buffer for both, which is what all call sites in tree do.
OpenCL does not define the result of reading from and writing to the same image object inside one kernel. Sampler reads and image writes are separate, non-coherent paths within a launch and work items are not ordered, so a work item can read a texel another work item has already replaced, or a stale copy of one. Being pointwise does not make the kernels safe here; it only means drivers that keep both paths in sync happen to hide the problem.
Where they do not, the conversion leaves scattered corrupt pixels behind and whatever consumes the buffer inherits them. Feeding such a buffer to the lanczos3 resample in 'scale into final size' turns single bad texels into non-finite output, reported by -d nan as:
[dev_pixelpipe] module `finalscale' outputs NaNs! [full]
Observed with rusticl on a Radeon RX 6700 XT; the same pipeline is clean under ROCm, which is why this went unnoticed.
Convert into a scratch image and copy the result into place instead.
Co-created with Claude Opus 5.0.