<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/135363>135363</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            [LLVM] Wrong assumption of `isKnownNonNullFromDominatingCondition`
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          DenisGZM
      </td>
    </tr>
</table>

<pre>
    Hello!

Right now `isKnownNonNullFromDominatingCondition` assumes that every target forbids division by zero and thus for every dominated usages of value we assume that it is not zero. It isn't true at least for CUDA

PTX ISA: "Division by zero yields an unspecified, machine-specific value"

Simple example here:
```
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void kernel2(int in1, int in2, int* out) {
 out[0] = (in1 / in2) - (in1 && in2);
}

int main() {
    int copy_out[1];
    int *d_out;

    cudaMalloc((void**)&d_out, 1 * sizeof(int));
    kernel2<<<1, 1>>>(1, 0,  d_out);
 cudaDeviceSynchronize();
    cudaMemcpy(copy_out, d_out, 1 * sizeof(int), cudaMemcpyDeviceToHost);

    printf("Device output: %d\n", copy_out[0]);
    return 0;
}
```
In EarlyCSEPass we "optimize" expression `(in1 / in2) - (in1 && in2)` -> `(in1 / in2) - (bool)(in1)`

Clang output: "Device output: -2"
NVCC output: "Device output: -1"

I think `isKnownNonNullFromDominatingCondition` interface should consider target architecture in case it wants to assume non zero value for divisor.
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJyUVUFv4zYT_TX0ZWCDGlmyffBBkeNvg28TLJrttujFoKWxxS5FGiTlrPPrC5JKnHTRbhcgbInDN2_ekHwSzsmjJlqz4oYVm4kYfGfsekNauv_9cT_Zm_ay_kBKGYYZ4xXj1S_y2HnQ5glYyaX7vzZP-sHoh0GprTX9xvRSCy_1sTa6lV4azUoOwrmhJwe-Ex7oTPYCXtgjeTgYu5etg1aepZNGw_4Cz2QNCN2C7wYXVoyQNiWnFgYnjuTAHOAs1EDwRCNFYpAepANtfEw1g7vwrhkuPHg7EAgPioSL7FD_uqmStk-ff4e7x4rlFTDEzd8rukhSrQOhYdDuRI08SGoZ1tCLppOapuNkk4piiCnto-xPioC-ifjfkSWWR8qSj4NXDHOpGzW0BCyvnW-lmXUsv_0u1Ayt-OfIzg7ay55eV_Bqtzsqsxdqt4OzkS18JatJIcOl1B6kzoKE9IjjI8MKzOAZroAtbhhPb8UNZ8UGWL6BCM6A4TahVjC9zpUMy3Ga5QHNFptUSmDphdQMl29yA0T6xpwuu8STsWKToC9BhlUbgylhCgTB90Ip08SEyyCPYRXHimGZEFhDqKoCJ5_JHJLuuGB15XhpSl6nEZuShR6mgcs4w8MPjHlf4KGMDZ1lQ48X3XTWaPlMSeKVINZKfXO6MFy-SsUaflRk_QaaWD6bD8Zd-UeCk5XaHyItpnVh006hZeE4Fy0rah3OZMh4bXXY0neFWvKD1cDfbd2bg3qn4VZYdakfbz8J58LVY4jm5GUfZSPQt5MlF69OgPzXo1JymLL89l8we2NU7MkyntrVeHN4VSuhj-_kft-CKaYL-fClrn-0NHu9u3fgO6m__pTZSe3JHkRD4DozqBYao51syb54nrBNJz01frAEUkMjHAXPehLaO_Dmxcu00cl5kskFs4o2aexs0q7zdpWvxITW2WI-zxcLLJaTbo00x-IgqGiLbM7nZTkvi4avcL_IOC4WfCLXyLHg8yzL8myeLWa83C_bvUBOPMfmULI5p15INVPq3M-MPU6kcwOts7zIy3yixJ6Ui18MRE1PEKOhYcVmYtcBNN0PR8fmXEnn3TWNl17FT83Hj1_ug5P8Zo0-JrGn0Lzg5z_R6Mlg1brz_uSCneKW4fYofTfsZ43pGW4D8fg3PVnzJzWe4TaW6xhuRz3nNf4VAAD__5K2KFY">