-
Notifications
You must be signed in to change notification settings - Fork 8
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
test @github-actions #2
Comments
@github-actions! |
@github-actions hello? |
welll... |
Collaborator
Author
wewe"\n- [ ] Run tests" |
qiao-bo
pushed a commit
that referenced
this issue
May 8, 2022
feisuzhu
pushed a commit
that referenced
this issue
Mar 29, 2023
…uilder for LLVM-CUDA backend (#5087) * [llvm] [aot] Add LLVM-CPU AOT tests * Refactored AOT test framework * Fixed minor issue * Enabled LLVM CPU-AOT for arm64 architecture * Added aot unit tests programming guide * [llvm] [aot] CUDA-AOT PR #2: Implemented AOT Module Loader for LLVM-CUDA backend * Fixed typo * Fixed minor issue * Refactored AOT test framework * [llvm] [aot] Add LLVM-CUDA AOT tests * Added cuda device availability check
feisuzhu
pushed a commit
that referenced
this issue
Mar 29, 2023
…uleBuilder to support Fields (#5120) * [aot] [llvm] Implemented FieldCacheData and refactored initialize_llvm_runtime_snodes() * Addressed compilation erros * [aot] [llvm] LLVM AOT Field #1: Adjust serialization/deserialization logics for FieldCacheData * [llvm] [aot] Added Field support for LLVM AOT * [aot] [llvm] LLVM AOT Field #2: Updated LLVM AOTModuleLoader & AOTModuleBuilder to support Fields * Fixed merge issues * Stopped abusing Program*
feisuzhu
pushed a commit
that referenced
this issue
Mar 29, 2023
…end (#5121) * [aot] [llvm] Implemented FieldCacheData and refactored initialize_llvm_runtime_snodes() * Addressed compilation erros * [aot] [llvm] LLVM AOT Field #1: Adjust serialization/deserialization logics for FieldCacheData * [llvm] [aot] Added Field support for LLVM AOT * [aot] [llvm] LLVM AOT Field #2: Updated LLVM AOTModuleLoader & AOTModuleBuilder to support Fields * [aot] [llvm] LLVM AOT Field #3: Added AOT tests for Fields running CPU backend * Added tests for activate/deactivate operations * Fixed merge issue
feisuzhu
pushed a commit
that referenced
this issue
Mar 29, 2023
… backend (#5124) * [aot] [llvm] Implemented FieldCacheData and refactored initialize_llvm_runtime_snodes() * Addressed compilation erros * [aot] [llvm] LLVM AOT Field #1: Adjust serialization/deserialization logics for FieldCacheData * [llvm] [aot] Added Field support for LLVM AOT * [aot] [llvm] LLVM AOT Field #2: Updated LLVM AOTModuleLoader & AOTModuleBuilder to support Fields * [aot] [llvm] LLVM AOT Field #3: Added AOT tests for Fields running CPU backend * Added tests for activate/deactivate operations * [aot] [llvm] LLVM AOT Field #4: Added AOT tests for Fields running CUDA backend * Fixed merge issues * Addressed naming issues
feisuzhu
pushed a commit
that referenced
this issue
Mar 29, 2023
Attempt #2 This PR got reverted due to failed release tests. ghstack-source-id: 7f4ee88e39858a8ae8d1d7012f3aa1bc08e0e037 Pull Request resolved: taichi-dev/taichi#6394
feisuzhu
pushed a commit
that referenced
this issue
Mar 29, 2023
Issue: taichi-dev/taichi#6434 ### Brief Summary 1. This is a special part of the Tacihi runtime module for the `AMDGPU` backend. Tacihi's runtime module uses `clang++` to generate `LLVM IR` is different in memory allocation differs from the cpu-generated `LLVM IR`. The following is an example. ``` C/C++ code void func(int *a, int *b) { *a = *b; } x86_64 backend LLVM IR define dso_local void @cpu_func(i32* %0, i32* %1) #2 { %3 = alloca i32*, align 8 %4 = alloca i32*, align 8 store i32* %0, i32** %3, align 8 store i32* %1, i32** %4, align 8 %5 = load i32*, i32** %4, align 8 %6 = load i32, i32* %5, align 4 %7 = load i32*, i32** %3, align 8 store i32 %6, i32* %7, align 4 ret void } __global__ function on AMDGPU define protected amdgpu_kernel void @global_func(i32 addrspace(1)* %0, i32 addrspace(1)* %1) #4 { %3 = alloca i32*, align 8, addrspace(5) %4 = alloca i32*, align 8, addrspace(5) %5 = alloca i32*, align 8, addrspace(5) %6 = alloca i32*, align 8, addrspace(5) %7 = addrspacecast i32* addrspace(5)* %3 to i32** %8 = addrspacecast i32* addrspace(5)* %4 to i32** %9 = addrspacecast i32* addrspace(5)* %5 to i32** %10 = addrspacecast i32* addrspace(5)* %6 to i32** %11 = addrspacecast i32 addrspace(1)* %0 to i32* store i32* %11, i32** %7, align 8 %12 = load i32*, i32** %7, align 8 %13 = addrspacecast i32 addrspace(1)* %1 to i32* store i32* %13, i32** %8, align 8 %14 = load i32*, i32** %8, align 8 store i32* %12, i32** %9, align 8 store i32* %14, i32** %10, align 8 %15 = load i32*, i32** %10, align 8 %16 = load i32, i32* %15, align 4 %17 = load i32*, i32** %9, align 8 store i32 %16, i32* %17, align 4 ret void } __device__ function on AMDGPU define hidden void @device_func(i32* %0, i32* %1) #2 { %3 = alloca i32*, align 8, addrspace(5) %4 = alloca i32*, align 8, addrspace(5) %5 = addrspacecast i32* addrspace(5)* %3 to i32** %6 = addrspacecast i32* addrspace(5)* %4 to i32** store i32* %0, i32** %5, align 8 store i32* %1, i32** %6, align 8 %7 = load i32*, i32** %6, align 8 %8 = load i32, i32* %7, align 4 %9 = load i32*, i32** %5, align 8 store i32 %8, i32* %9, align 4 ret void } ``` 2. There are some differences in the place about `allocainst`, specifically about addrspace (for `AMDGPU`, [this](https://llvm.org/docs/AMDGPUUsage.html#address-spaces) will be helpful). I have not found documentation describing how to write the correct `LLVM IR` on `AMDGPU`, through my observation of the `LLVM IR` generated by `clang++/hipcc`. We need to deal with the arguments of the `__global__` function and the `allocainst` (including specifying the addrspace of `allocainst` and performing addrspace-cast) while for the `__device__` function we do not need to deal with the arguments of the function. Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
feisuzhu
added a commit
that referenced
this issue
Mar 29, 2023
Related Issue: #7140 ### Brief Summary On macOS, when test worker hard fails (abort, EXC_BAD_ACCESS, etc.), backward_cpp's signal handler will re-raise the signal and catch it afterwards, make it an infinite loop, at the moment the offending process can't be terminated easily (except a SIGKILL), eat CPU cycles and blocks test runner. ``` (lldb) bt * thread #1, queue = 'com.apple.main-thread', stop reason = signal SIGABRT * frame #0: 0x00000001a04f0e28 libsystem_kernel.dylib`__pthread_kill + 8 frame #1: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #2: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #3: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #4: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #5: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #6: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #7: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #8: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #9: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #10: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #11: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #12: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #13: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #14: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #15: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #16: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #17: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #18: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #19: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #20: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #21: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #22: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #23: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #24: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #25: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #26: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #27: 0x00000001283a0848 taichi_python.cpython-38-darwin.so`backward::SignalHandling::sig_handler(int, __siginfo*, void*) + 28 frame #28: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #29: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #30: 0x00000001a0402e10 libsystem_c.dylib`raise + 32 frame #31: 0x00000001a056ec44 libsystem_platform.dylib`_sigtramp + 56 frame #32: 0x00000001a052343c libsystem_pthread.dylib`pthread_kill + 292 frame #33: 0x00000001a046b454 libsystem_c.dylib`abort + 124 frame #34: 0x0000000100194fc0 python`os_abort + 12 frame #35: 0x00000001000758a8 python`cfunction_vectorcall_NOARGS + 324 frame #36: 0x00000001001140f0 python`call_function + 460 frame #37: 0x000000010011086c python`_PyEval_EvalFrameDefault + 27176 frame #38: 0x00000001000287e4 python`function_code_fastcall + 128 frame #39: 0x0000000100028008 python`PyVectorcall_Call + 120 frame #40: 0x0000000100110b20 python`_PyEval_EvalFrameDefault + 27868 frame #41: 0x000000010010982c python`_PyEval_EvalCodeWithName + 3008 frame #42: 0x0000000100028948 python`_PyFunction_Vectorcall + 208 frame #43: 0x0000000100028008 python`PyVectorcall_Call + 120 ```
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
test @github-actions???
The text was updated successfully, but these errors were encountered: