Add review_requested.yml #4
Closed
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
test comment.yml
This reverts commit 47f85c9.
archibate
requested review from
k-ye and
taichi-gardener
and removed request for
taichi-gardener and
k-ye
feisuzhu
pushed a commit
that referenced
this pull request
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 pull request
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 pull request
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
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
No description provided.