Atomic add on SharedArray vector causes LLVM assertion

[oliver-batchelor]

Using an atomic add on a SharedArray vector causes the assertion below.
(May be related to #8435?)

> python bad_shared.py
[Taichi] version 1.7.0, llvm 15.0.4, commit 6a30b8dc, linux, python 3.10.11
[Taichi] Starting on arch=cuda
python: /home/ailzhang/github/llvm-project/llvm/lib/IR/Instructions.cpp:1647: void llvm::AtomicRMWInst::Init(llvm::AtomicRMWInst::BinOp, llvm::Value *, llvm::Value *, llvm::Align, llvm::AtomicOrdering, SyncScope::ID): Assertion `cast<PointerType>(getOperand(0)->getType()) ->isOpaqueOrPointeeTypeMatches(getOperand(1)->getType()) && "Ptr must be a pointer to Val type!"' failed.
[1]    28401 IOT instruction (core dumped)  python bad_shared.py

import taichi as ti

@ti.kernel
def foo():

  ti.loop_config(block_dim=64)
  for x in range(1000):
    shared = ti.simt.block.SharedArray((64, ), dtype=ti.math.vec3)
    shared[x % 64] = ti.math.vec3(0)
    ti.simt.block.sync()
    
    for i in range(64):
      ti.atomic_add(shared[i], ti.math.vec3(1))

ti.init(arch=ti.cuda)
foo()

[bobcao3]

Does float work in the shared array?

[oliver-batchelor]

Yes, and also updating the components one by one with atomicAdd also works.

…

On Fri, Dec 29, 2023 at 7:01 AM Bob Cao ***@***.***> wrote: Does float work in the shared array? — Reply to this email directly, view it on GitHub <#8456 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AAITRZKXLKQI24BS7W5ZHZTYLWXXTAVCNFSM6AAAAABBFP6GWCVHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHMYTQNZRGM4DINZRGQ> . You are receiving this because you authored the thread.Message ID: ***@***.***>