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

    <tr>
        <th>Summary</th>
        <td>
            Failing assertion in AMDGPUAttributor with ptrtoint casts and AS 3
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            backend:AMDGPU,
            crash-on-valid
      </td>
    </tr>

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

    <tr>
      <th>Reporter</th>
      <td>
          ritter-x2a
      </td>
    </tr>
</table>

<pre>
    I observe a failing assertion in the AMDGPUAttributor in code with ptrtoint casts and address space 3. It occurs on trunk since commit 41ed16c3b3362e51b7063eaef6461ab704c1ec7a by @jwanggit86.

Reproducer:
```
target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"
target triple = "amdgcn-amd-amdhsa"

@buf_shared = internal addrspace(3) global [2080 x i8] undef, align 16

define protected amdgpu_kernel void @foo(ptr addrspace(1) nocapture noundef writeonly initializes((0, 1)) %res.coerce) local_unnamed_addr {
entry:
  %conv.i = and i32 trunc (i64 sub (i64 16, i64 ptrtoint (ptr addrspacecast (ptr addrspace(3) @buf_shared to ptr) to i64)) to i32), 15
  %add.ptr = getelementptr inbounds nuw i8, ptr addrspace(3) @buf_shared, i32 %conv.i
  %0 = load i8, ptr addrspace(3) %add.ptr, align 1
  store i8 %0, ptr addrspace(1) %res.coerce, align 1
  ret void
}
```

`opt -mcpu=gfx1030 --amdgpu-attributor frame.ll` with the above as `frame.ll` yields:

```
opt: /home/faritter/projects/ritter-x2a-fork/llvm-project/llvm/include/llvm/Support/Casting.h:578: decltype(auto) llvm::cast(From*) [with To = llvm::PointerType; From = llvm::Type]: Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: ../../build/bin/opt -mcpu=gfx1030 --amdgpu-attributor frame.ll
1.      Running pass 'AMDGPU Attributor' on module 'frame.ll'.
 #0 0x00005d350dd6d050 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (../../build/bin/opt+0x531f050)
 #1 0x00005d350dd6a46f llvm::sys::RunSignalHandlers() (../../build/bin/opt+0x531c46f)
 #2 0x00005d350dd6a5c5 SignalHandler(int) Signals.cpp:0:0
 #3 0x000078045a842520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x000078045a8969fc __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
 #5 0x000078045a8969fc __pthread_kill_internal ./nptl/pthread_kill.c:78:10
 #6 0x000078045a8969fc pthread_kill ./nptl/pthread_kill.c:89:10
 #7 0x000078045a842476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #8 0x000078045a8287f3 abort ./stdlib/abort.c:81:7
 #9 0x000078045a82871b _nl_load_domain ./intl/loadmsgcat.c:1177:9
#10 0x000078045a839e96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
#11 0x00005d3509d95b2b (anonymous namespace)::AAAMDAttributesFunction::needFlatScratchInit(llvm::Attributor&) AMDGPUAttributor.cpp:0:0
#12 0x00005d3509da0aeb (anonymous namespace)::AAAMDAttributesFunction::updateImpl(llvm::Attributor&) AMDGPUAttributor.cpp:0:0
#13 0x00005d350c2c03ea llvm::AbstractAttribute::update(llvm::Attributor&) (../../build/bin/opt+0x38723ea)
#14 0x00005d350c2d319d llvm::Attributor::updateAA(llvm::AbstractAttribute&) (../../build/bin/opt+0x388519d)
#15 0x00005d3509d9fb96 (anonymous namespace)::AAAMDAttributes const* llvm::Attributor::getOrCreateAAFor<(anonymous namespace)::AAAMDAttributes>(llvm::IRPosition, llvm::AbstractAttribute const*, llvm::DepClassTy, bool, bool) (.constprop.0) AMDGPUAttributor.cpp:0:0
#16 0x00005d3509da213b (anonymous namespace)::runImpl(llvm::Module&, llvm::AnalysisGetter&, llvm::TargetMachine&, llvm::AMDGPUAttributorOptions) (.constprop.0) AMDGPUAttributor.cpp:0:0
#17 0x00005d3509da2a3e (anonymous namespace)::AMDGPUAttributorLegacy::runOnModule(llvm::Module&) AMDGPUAttributor.cpp:0:0
#18 0x00005d350db050a9 llvm::legacy::PassManagerImpl::run(llvm::Module&) (../../build/bin/opt+0x50b70a9)
#19 0x00005d35095b6d12 optMain (../../build/bin/opt+0xb68d12)
#20 0x000078045a829d90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#21 0x000078045a829e40 call_init ./csu/../csu/libc-start.c:128:20
#22 0x000078045a829e40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#23 0x00005d35095ac855 _start (../../build/bin/opt+0xb5e855)
Aborted (core dumped)
```

Used cmake options (probably not minimal):
```
-DCMAKE_BUILD_TYPE:STRING=Release
-DCMAKE_CXX_COMPILER_LAUNCHER:STRING=ccache
"-DLLVM_ENABLE_PROJECTS:STRING=clang;lld;clang-tools-extra"
"-DLLVM_ENABLE_RUNTIMES:STRING=compiler-rt;openmp"
-DLLVM_ENABLE_ASSERTIONS:BOOL=TRUE
-DLLVM_FORCE_ENABLE_STATS:BOOL=TRUE
-DLLVM_USE_SPLIT_DWARF:BOOL=TRUE
-DLLVM_ENABLE_DUMP:BOOL=TRUE
```

I reduced the above IR from this HIP code:
```
// clang -xhip --offload-arch=gfx1030 -isystem /opt/rocm/include --driver-mode=g++ -O3 ./frame.hip

#include "hip/hip_runtime.h"

#define ALIGNMENT_ZERO_BITS 4
#define NUM_MOVE_THREADS 64
#define BYTES_PER_THREAD 32
#define ALIGNMENT_MASK ((1u << ALIGNMENT_ZERO_BITS) - 1)
#define BUFFER_SIZE (NUM_MOVE_THREADS * BYTES_PER_THREAD)
#define ADD_ALIGN_SLACK(e) ((e) + 2 * (ALIGNMENT_MASK + 1))
#define BUFFER_ALLOC_SIZE ADD_ALIGN_SLACK(BUFFER_SIZE)

__host__ __device__
uint64_t compute_alignment_offset(uint8_t *orig_ptr) {
  uint64_t ptr = (uint64_t) orig_ptr;
  uint64_t alignment_bits = (ptr & ALIGNMENT_MASK);
  return ((ALIGNMENT_MASK - alignment_bits) + 1) & ALIGNMENT_MASK;
}

__global__
void MoveKernelThroughput(uint8_t * res) {
  __shared__ uint8_t buf_shared[BUFFER_ALLOC_SIZE];
  uint8_t *buf_adjusted = buf_shared + compute_alignment_offset(buf_shared);
  *res = *buf_adjusted;
}

```
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJykWVtz46iz_zTkhZJLAt38kAf5NuMzcZKynT3__76oEMI2GxlUgDzx-fSnQPLdk83sTk1sCbp_3fQNaBOt-Vow9giiAYhGD6QxG6keFTeGKe8DkYdClvvHKZSFZmrHIIErwisu1pBozZThUkAuoNkwmM1G317fMmMULxojlR2nsmTwJzcbWBtlJBcGUqKNhkSUkJSlYlpDXRPKIO7BqYGS0kZpKAU0qhHvUHNBGaRyu-UGhgErg5jiAuMYsSgoEj_GjLBVHMYBKRI_pAGjCYHFHoLQ_-snEes1N2ncA34G_GzOaiXLhjIFsH0Hsd_99zND1JoZWBJDKrKXjYEAjyBAiHk1wFkcug-vDs5fEMAZbj-8Gp-_hOdk0flMfP6SAJwFsQ9whqL4-GlnUjuDuk-v7tuHProl5AcpuyB2YN4Ohe3U7iBmF6YdTT_uAHdHMG93RPR2UWBHowB5u8B3OPbL2yHfQdgvTzjYOPQWGHlZ5H0LPMEBzuxaLFEfIHQyqFG8rtjBmGRbrqnwyLa0fxtNWlr7P_SLZpXrDVGsdORcGKYEqVycuBgBKMUA9eG6kgWpIIgGyE99-AF5CqIRbETJVgANIan4WsAgbpFLtuKCwVpJw6hhJbRK1E3-zpRgFdxJXtpoWUkJUFobdSEvsPKEpKQ2jWJQSCcE_lTcMCmqPeSCG04q_n9MA5QClPpWA8tmOQGKFNM9KpmycH1YSUqqvBGCbFmZW0kQJAPgZ0wYtW_DElo2KsWux50hbK5wjFxGUAhQyuMQ6qY4PAaxFWmfjkl2vRKbdDeDB3Nemt5IC2PHjbSg3UrsC0buZQiD6KAmKcueBbV6rplhFdsyYewIF4U1loai-WkdhIbw76W7hWB0MsBBju8kVJKUn2IdFToLAwehjVQM8tRh3eMP7rjrEkIx44LFRlUyuiof7ausDfS2tG4AHq1XH4GPfeh5bbx55FQZV4psWa-qQOy31dHWT1JIW2A1BLF_TrDnrCp1V7EupcraAJxBgCYbuWUATVakLd0ATWol_2LUaIAmp3LuraR6B2hSVbut11F0rwBNuKBVU7LTwKKpa6ksxZBow8W6t7HlIbFpDktGK7OvrfFIY6SLbsuGM4AzG3AApRMltwBlzrbRwC11KVtXHklfpcv0pYXCA2hZrijcVDSyQrPjrgNin2sC8HApAR4DlP5BqtaHMUCxLTZOBzxc7tt5O0nUurEBCuUKckHltiaGFxWD7UICW45Q4vY4Vtpd4_VpnC3GNtvsBkRg0ayhYtYoNiM2xtTOM2gC0GTNzaYpetQu-WDBK0NzrRtbKCZtUrf2dt6niugNLAh9N4pQZmUvDKHvsGy2det8vwfdv1cl14psj4uxGsBeD6CJ-ygaXpX2mwuAJr8dkn4WdHLmjRB2p6-J1hCgpN3g4WmHt6aSAm5l2dgSj5IjCErsAiBA2If-h-_7flTiyC_LuPQj_8y1eq-7KFBcGLfipWoz8kSkyM9caqMY2TrvDu3e0Po6_WTdAA38jwgHK9-mfL9TKLhSiITx6p5C80Ys-FqQ6jsRZcWU7kLoazJpGK9OMtG1zIhG8ALdlvJ2Te2w7tHaut13fy0K7lCS1A8jkoYoQj50Wk0qXgA0-UjjPA69iovmw1uLpp2gPS17sVPMsZzUCi8B-3F_RWGe12ajGCnzd15VOd_WbU0nLuvsokVtKltgzsh6FOAstOeFJO7Ao6-AH7b4z2BdtQkORojvwZ7zfIqV9s-xkmuDhkkM19o5wMG0jwdn670uWW2zt5aaf9jCSrhmDhjZ089h6eklLkqTFbblXZkW1ZStv9xQq1fgDlAte_-GPShgLqrcbn95KbeEt47gwq3SDm_1mpIWKwiSxB3D_MyGu3-JhvusH_9e1DiWNmos4EUC9ct-VCB3FCFCiv1WNhraw023r_bbZMqybDY6FA6mJ42gNpzaScFYOamIWVBFDN1MBTcX6X9ecKweNzeNq1yxSqJLJYlP2L9TsqlLYth0W1f_Wjd8rhtF1MeMnFWgrNB2CzBHTc4V-Fz4F2oTThOEGTm5M7zUpsRBv4R3ZZypkWWXityo_Bv6pFHQL0_6RFfhtSraeP0Nz0EqhT1-ZJ-sY83Mixoq5hYzsaPD3xPSHipOAqbzV6m5Cxg0_MydR-0u6UasHlZE6-XejhdSVqfv1pKOr1ay7vlfjLT4KgtQgP82C1QjbqJ85nb4bu89W5og1V5z_Y21Z86r6aW7Ac4I3XBxh_lK_Zfa2k7_88Um14slmP1t4FzBPrE1ofujIV7EYeX3jfElvdKLzb_wI5-cn5SrM5GvROsZEWTNlPPBQY9fy__KacQvEp_0TxnWvzBUVMRlgKCszczuK18ALOK0DNARD13tMKhf9n2Y53YTybUhyuSUVFV-3Laobm631G7HvsfUXjrcKSDuRAbXIlnoQ8dg7-PXUtpHC-056HabRG1Ho0NE9xAvFnFP_7vIOLHRFXXA-NLahKZRBFvILxk7YmkUtcbO7HmBlZaN2uusvRqwQ-m8vIy-aVZCuiXvzHrWJpa7_itZkKLaQyEN3HLBt-7OdNsN80bDWfZjnA_epk-jfPnf1zHA2WI5nz5_A3g0ZxUjmp3RDf_zn3z4MnudPo3n-VP29jz8Pp6fs1BK6IY5kyBv9PT0xywfP2eDp3H-On_5n_FwubigrohYAzyoqhLggXvzjJSV9tiHUYem0TXS_O15OZ2NL5HktuYVU54yAA9kzcS2btkvebPFYjxfTl-eLffg5eUJ4NFy_jY-EU5e5sPxgXyxzJa_pHxbjPPF69N0mY_-N5tPfkXWQY3eZq83JFfenELFyoay8qxRMJ3Dlb0qmw3X8Pv01fVabz3ZXk2hsyH0Pja8hp4nVyt7aPSIopvzeyHXe23YFh7ib6IkPWsMQM8rFd8x5W2drNEaoAFAA-i9YJcZ7f1vw-uuVYHwgRMgZIfRZMPrXDXCcEt46v4h3LXpsqfpt-fZ-HmZ_zmev-SD6XIBwwuK57dZPnv5Y5wvv8_H2WgB48v5wX-X40X-Op53BBCjX0iYZYsf7VE4DRpoTwB4eE8BW2i9tqd3IehtMhnP88X0z7FFudHLHkCulbnGyEaj3EnMF0_Z8AdAKevK-ulxAJHDAii90X1waDXe0yx7enoZtvrdyjnTvmP3szzfSG3yHOZ5yXacsjwHftZwYeIwN9AmU2NY7rpi9k6Yy9VKM3tatzRpbitaJhVf510Lse1tQniEOPQKOxY7ZumOPPiK_iSq4EYfWB0Kiq9c6SrZ4NCra5TozHhlNO8K9GDkrgN4g-ogu46fNVHbfXaWcb3jmdyxH66XvNwo2aw3dXNlEaiYPrdG3nU78xweqM5aoNHgxn2u-XWyTAdreUj5V6NN1zE_b6CjwWfuOu-4Ho0GUKbYwcYX6Fc2ONaXh_IRl33cJw_sMUhwGEbID8KHzSPB_VWQJBGmCfOjNEU0oX6QJgkJ-n6cRA_8EfkoDFCQBHGAcdgrS9SnSRRGSRKliKQg9NmW8KpnDz49qdYPrnf2GCAfRfFDRQpWafezFUIFoe9MlMfznOviDQFCrqnmSeHtSMVLOxyNHtSja8oVzVqD0K-4NvokxHBTscfJvd-4bn7f-tXvWtkC4odGVY__vD3YrXL3iP4_AAD__7T7Mo0">