[lang] Add popcnt to llvm intrinsic support

[lgyStoic]

[pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

fix: add test api for popcnt

Issue: #7541

Brief Summary

🤖 Generated by Copilot at 9893837

Added a new popcnt function to count the number of bits set to 1 in an integer, both for Python and Taichi expressions. Implemented the function in the taichi.lang.ops and taichi.math modules, and the CUDA and LLVM backends. Added tests and documentation for the function.

Walkthrough

🤖 Generated by Copilot at 9893837

• Add a new function popcnt to the taichi.lang.ops module that returns the number of bits set to 1 in an integer (link, link)
• Import and wrap the popcnt function in the taichi.math.mathimpl module to support both Taichi expressions and Python scalars (link, link, link)
• Add the popcnt operation to the UnaryOpType enum class and the expression macros in the taichi.inc.unary_op.inc.h and taichi.ir.expression_ops.h files (link, link)
• Add the popcnt operation to the UNARY_STD functions in the taichi.codegen.cuda.codegen_cuda.cpp and taichi.codegen.llvm.codegen_llvm.cpp files, which handle the code generation for unary operations on the CUDA and LLVM backends (link, link)
• Add a temporary workaround for creating a UnaryOpStmt with the popcnt operation using a BinaryOpStmt with the min operation in the IRBuilder class in the taichi.ir.ir_builder.cpp and taichi.ir.ir_builder.h files (link, link)
• Export the popcnt function to the Python module taichi.lang.expr using the DEFINE_EXPRESSION_OP macro in the taichi.python.export_lang.cpp file (link)
• Add the popcnt function to the list of math functions that are tested for consistency between Taichi expressions and Python scalars in the tests.python.test_api.py file (link)

[netlify]

✅ Deploy Preview for docsite-preview ready!

Name	Link
🔨 Latest commit	2a24e82
🔍 Latest deploy log	https://app.netlify.com/sites/docsite-preview/deploys/643d5e38fc57910008eef930
😎 Deploy Preview	https://deploy-preview-7772--docsite-preview.netlify.app
📱 Preview on mobile	Toggle QR Code... Use your smartphone camera to open QR code link.

---

To edit notification comments on pull requests, go to your Netlify site settings.

[lgyStoic]

@ailzhang would you please have a look what is going on, if I use

import taichi as ti
ti.init(arch=ti.cuda, print_kernel_llvm_ir_optimized=True)
@ti.kernel
def test():
    n = ti.uint64(10000)
    ti.math.max(10, n)
    pt = ti.math.popcnt(n)
    print(pt)

test()

using uint64 as popcnt type, then I pass to cuda backend for compute, then assert like the picture mentioned.

[ailzhang]

@lgyStoic the error message indicates that you're passing a pointer to u64 to a function requires a pointer to i64. you'll need to use builder->CreateBitCast to cast it to i64 before sending to the function.

[lgyStoic]

@lgyStoic the error message indicates that you're passing a pointer to u64 to a function requires a pointer to i64. you'll need to use builder->CreateBitCast to cast it to i64 before sending to the function.

At last, found beause return type of popcll is i32, but in statement return type pointer is to u64 point, so the print function was crash.... at first, I was thinking there is some problem on call 'popcll' function..

[lgyStoic]

/rebase

[turbo0628]

LGTM!