Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

LLVM Fatal Error: Copy one register into another with a different width #8086

Closed
AmesingFlank opened this issue May 26, 2023 · 4 comments
Closed
Assignees

Comments

@AmesingFlank
Copy link
Collaborator

AmesingFlank commented May 26, 2023

repro:

import taichi as ti

ti.init(ti.cuda)

@ti.dataclass
class S:
    i: ti.i16 = 0
    b: bool = False

sf = S.field(shape=(10,  2))

@ti.kernel
def k() -> S:
    ss = [sf[0,0], sf[0,1]]
    return ss[0]

k()

this throws:

Traceback (most recent call last):
  File "C:\Users\ldfra\Code\taichi\test_llvm_throw.py", line 19, in <module>
    k()
  File "c:\users\ldfra\code\taichi\python\taichi\lang\kernel_impl.py", line 936, in wrapped
    return primal(*args, **kwargs)
  File "c:\users\ldfra\code\taichi\python\taichi\lang\kernel_impl.py", line 868, in __call__
    return self.launch_kernel(kernel_cpp, *args)
  File "c:\users\ldfra\code\taichi\python\taichi\lang\kernel_impl.py", line 801, in launch_kernel
    raise e from None
  File "c:\users\ldfra\code\taichi\python\taichi\lang\kernel_impl.py", line 798, in launch_kernel
    prog.launch_kernel(compiled_kernel_data, launch_ctx)
RuntimeError: [llvm_context.cpp:taichi::lang::TaichiLLVMContext::{ctor}::<lambda_ea41ae353cd42ca584262925730f46b7>::operator ()@80] LLVM Fatal Error: Copy one register into another with a different width

The problem looks to be the b:bool = False member, because removing it or changing its type removes the error.

@ailzhang
Copy link
Contributor

ailzhang commented Jun 1, 2023

cc: @listerily , maybe this is related to the boolean type?

@listerily
Copy link
Contributor

cc: @listerily , maybe this is related to the boolean type?

Seems yes. Maybe there is some special code generation rules related to field? I would look into it later.

@jim19930609 jim19930609 moved this from Untriaged to Todo in Taichi Lang Jun 2, 2023
@listerily
Copy link
Contributor

This issue is caused by TaskCodeGenCUDA::create_intrinsic_load. I'll keeping on working to solve that.

listerily added a commit that referenced this issue Jun 5, 2023
Issue: #8086 

### Brief Summary

This pull request adds support for bool types in
`TaskCodeGenCUDA::create_intrinsic_load` which solves bit width issue on
CUDA runtime. The problem was that `nvvm_ldg_global_i` does not support
1 bit integer. So we cast pointers and values to `i8` to solve this
issue.

### Walkthrough

+ Added value and pointer casting in
`TaskCodeGenCUDA::create_intrinsic_load`
+ Added test case

---------

Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
@listerily
Copy link
Contributor

Hi! Your issue has been solved and it has been merged into master. Additionally, default values are not supported in taichi dataclasses. Please remove default values in S like this:

@ti.dataclass
class S:
    i: ti.i16
    b: bool

@github-project-automation github-project-automation bot moved this from Todo to Done in Taichi Lang Jun 5, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
Status: Done
Development

No branches or pull requests

3 participants