The magic numbers from the nextBaseAddress are incorrectly compared with constant values

[AlexDemydenko]

constant values. They are usually simplified as unequal values.

[google-cla]

Thanks for your pull request! It looks like this may be your first contribution to a Google open source project. Before we can look at your pull request, you'll need to sign a Contributor License Agreement (CLA).

View this failed invocation of the CLA check for more information.

For the most up to date status, view the checks section at the bottom of the pull request.

[rjodinchr]

Could you provide a small test exercising the issue?

[AlexDemydenko]

CL code

kernel void test_kernel(global float *src, global long *dst)
{
    uint tid = get_global_id(0);
    dst[tid] = (long)(src != 0);
}

SPIRV below is missing:

................
         %40 = OpPtrNotEqual %bool %30 %38
         %43 = OpSelect %ulong %40 %ulong_1 %ulong_0
.............

The pass just will remove it, because the src can't be equal to nullptr.

[rjodinchr]

Can you add a unit test with that to the PR please?

[rjodinchr]

Can you sign the CLA please? We won't be able to merge this PR without it. Thanks

[AlexDemydenko]

Could you please merge this PR?

[rjodinchr]

Your branch is 53 commits behind google/main. Please rebase on top-of-tree in order to have it tested properly by the bots.
Thank you

[AlexDemydenko]

done

[rjodinchr]

Can you explain why it is incorrect to compare the magic numbers with constant values? This is not obvious to me right now, and this pass has been designed to do exactly that. So this change is basically trying to get rid of that pass.

Maybe you just want to call clspv with --hack-logical-ptrtoint=0?

[AlexDemydenko]

The LogicalPointerToIntPass generates the base-addresses for each argument in some virtual/compiler memory space.
arg1 = 0x10000000ULL
arg2 = 0x20000000ULL
arg3 = 0x30000000ULL

This is not a real address, but it is good enough to compare (==, !=) it with other generated pointers (It also has problems, but that's not what this commit about).
But it has no sense to compare generated-address to any pointer from runtime/GPU memory space.
A constant in the kernel contains a value from runtime/GPU memory space.

If you compare any generated-pointer to nullptr it will always be false in terms of logic of the LogicalPointerToIntPass, but tis is not true.

[rjodinchr]

Do you mean that your issue is with a comparison returning true because a constant matches the virtual address LogicalPointerToIntPass has generated for an object (while it should return false)?

[AlexDemydenko]

It is return true because a constant NOT matches the virtual address LogicalPointerToIntPass has generated.
dst[tid] = (long)(src != 0);
as result replace OpPtrNotEqual instruction to a "True" value.

[rjodinchr]

Is this about the OpenCL-CTS test_api null_buffer_arg?
If it is, note that you definitively need physical_addressing enabled to pass conformance.

[AlexDemydenko]

Yes, this is about this test. And we want to pass this test. This test is passed with this change.

[rjodinchr]

While this change is fixing the test on your end, I'm not sure this is the way to do it properly for clspv.
I believe this test needs physical_addressing enabled to pass properly.

[AlexDemydenko]

This pass has to handle Logical pointers and it doesn't require the physical-storage-buffers flag. What is "the way" to solve this issue at your opinion?

Before I will get response from OCL team. I can try to do another solution, but I need to understand what are you thinking about this?

[rjodinchr]

I think that OpenCL test (test_api null_buffer_arg) requires physical-storage-buffers to pass properly with the Vulkan-Validation-Layers enabled.

You will note that we have this back and forth discussion because you are not sharing enough to let us understand what is happening on your side (you could have mentioned the OpenCL-CTS directly in your PR for example).

What I don't understand right now is how you have set up Vulkan so that OpPtrNotEqual works without physical-storage-buffers.