search

CHIP-SPV / chipStar

chipStar is a tool for compiling and running HIP/CUDA on SPIR-V via OpenCL or Level Zero APIs.
Other
219 stars 32 forks source link

Compiler crash on LLVM Texture pass #177

Open pvelesko opened 2 years ago

pvelesko commented 2 years ago

Uncovered while working on CHIP-SPV integration into HIP. Kernel:

// read from a texture using normalized coordinates
constexpr size_t ChannelToRead = 1;
template <typename T>
__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
                                bool textureGather) {
  #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
  // Calculate normalized texture coordinates
  const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  const float u = x / (float)width;

  // Read from texture and write to global memory
  if (height == 0) {
    output[x] = tex1D<T>(texObj, u);
  } else {
    const float v = y / (float)height;
    if (textureGather) {
      // tex2Dgather not supported on __gfx90a__
      #if !defined(__gfx90a__) && !(defined(__HIP_PLATFORM_SPIRV__))
      output[y * width + x] = tex2Dgather<T>(texObj, u, v, ChannelToRead);
      #else
      #warning("tex2Dgather not supported on gfx90a");
      #endif
    } else {
      output[y * width + x] = tex2D<T>(texObj, u, v);
    }
  }
  #endif
}

Error:

Consolidate compiler generated dependencies of target MemoryTest
[ 51%] Building CXX object catch/catch_tests/unit/memory/CMakeFiles/MemoryTest.dir/hipMallocArray.cc.o
In file included from /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/unit/memory/hipMallocArray.cc:31:
/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/unit/memory/hipArrayCommon.hh:110:8: warning: ("tex2Dgather not supported on gfx90a"); [-W#warnings]
      #warning("tex2Dgather not supported on gfx90a");
       ^
1 warning generated when compiling for .
Don't know how to lower this texture use case.
UNREACHABLE executed at /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:380!
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.  Program arguments: /gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt /tmp/hipMallocArray-generic-link-6789ea.bc -load-pass-plugin /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build/lib/libLLVMHipSpvPasses.so -passes=hip-post-link-passes -o /tmp/hipMallocArray-generic-lower-0e3839.bc
 #0 0x0000000002a729b3 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x2a729b3)
 #1 0x0000000002a73097 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f800b26c8c0 __restore_rt (/lib64/libpthread.so.0+0x168c0)
 #3 0x00007f8009d23cdb raise (/lib64/libc.so.6+0x4acdb)
 #4 0x00007f8009d25375 abort (/lib64/libc.so.6+0x4c375)
 #5 0x00000000029dea7f (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x29dea7f)
 #6 0x00007f8009a84dec (anonymous namespace)::lowerTextureObjectUses(llvm::Function*, std::vector<(anonymous namespace)::TextureUseGroup, std::allocator<(anonymous namespace)::TextureUseGroup>> const&) /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:332:32
 #7 0x00007f8009a83f4d (anonymous namespace)::lowerTextureFunctions(llvm::Module&) /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:482:13
 #8 0x00007f8009a83a5c HipTextureLoweringPass::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:492:10
 #9 0x00007f8009a4d4a4 llvm::detail::PassModel<llvm::Module, HipTextureLoweringPass, llvm::PreservedAnalyses, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) /gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/include/llvm/IR/PassManagerInternal.h:88:17
#10 0x000000000228812c llvm::PassManager<llvm::Module, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x228812c)
#11 0x000000000078c1a3 llvm::runPassPipeline(llvm::StringRef, llvm::Module&, llvm::TargetMachine*, llvm::TargetLibraryInfoImpl*, llvm::ToolOutputFile*, llvm::ToolOutputFile*, llvm::ToolOutputFile*, llvm::StringRef, llvm::ArrayRef<llvm::StringRef>, llvm::ArrayRef<llvm::PassPlugin>, llvm::opt_tool::OutputKind, llvm::opt_tool::VerifierKind, bool, bool, bool, bool, bool, bool) (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x78c1a3)
#12 0x000000000079cc13 main (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x79cc13)
#13 0x00007f8009d0e2bd __libc_start_main (/lib64/libc.so.6+0x352bd)
#14 0x0000000000786a2a _start /home/abuild/rpmbuild/BUILD/glibc-2.31/csu/../sysdeps/x86_64/start.S:122:0
clang-15: error: unable to execute command: Aborted (core dumped)
clang-15: error: hipspv-link command failed due to signal (use -v to see invocation)
clang version 15.0.0 (https://github.com/llvm/llvm-project.git 66fa2847a775dda27ddcac3832769441727db42f)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin
clang-15: note: diagnostic msg: Error generating preprocessed source(s).

failed to execute:/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/clang++ -D__HIP_PLATFORM_SPIRV__= -x hip --target=x86_64-linux-gnu -Xclang -no-opaque-pointers --offload=spirv64 -nohipwrapperinc --hip-path=/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build -isystem "/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build/include"  -DKERNELS_PATH="/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/kernels/" -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/external/Catch2 -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/./include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/./kernels -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build/include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/external/picojson -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/include/cuspv -Wno-duplicate-decl-specifier -Wno-tautological-constant-compare -Wno-c++20-extensions -Wno-unused-result -Wno-delete-abstract-non-virtual-dtor -Wno-deprecated-declarations -Wunused-command-line-argument -gdwarf-4 --std=c++17 -g -fPIE -Wno-format-extra-args -Wall -std=c++14 -pthread -std=c++17 -MD -MT catch/catch_tests/unit/memory/CMakeFiles/MemoryTest.dir/hipMallocArray.cc.o -MF CMakeFiles/MemoryTest.dir/hipMallocArray.cc.o.d -o CMakeFiles/MemoryTest.dir/hipMallocArray.cc.o -c /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/unit/memory/hipMallocArray.cc
pvelesko commented 2 years ago

one more:

Consolidate compiler generated dependencies of target TextureTest
[ 89%] Building CXX object catch/catch_tests/unit/texture/CMakeFiles/TextureTest.dir/hipTextureRef2D.cc.o
Don't know how to lower this texture use case.
UNREACHABLE executed at /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:380!
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.  Program arguments: /gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt /tmp/hipTextureRef2D-generic-link-f58c4d.bc -load-pass-plugin /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build/lib/libLLVMHipSpvPasses.so -passes=hip-post-link-passes -o /tmp/hipTextureRef2D-generic-lower-173736.bc
 #0 0x0000000002a729b3 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x2a729b3)
 #1 0x0000000002a73097 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f8d71d628c0 __restore_rt (/lib64/libpthread.so.0+0x168c0)
 #3 0x00007f8d70819cdb raise (/lib64/libc.so.6+0x4acdb)
 #4 0x00007f8d7081b375 abort (/lib64/libc.so.6+0x4c375)
 #5 0x00000000029dea7f (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x29dea7f)
 #6 0x00007f8d7057adec (anonymous namespace)::lowerTextureObjectUses(llvm::Function*, std::vector<(anonymous namespace)::TextureUseGroup, std::allocator<(anonymous namespace)::TextureUseGroup>> const&) /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:332:32
 #7 0x00007f8d70579f4d (anonymous namespace)::lowerTextureFunctions(llvm::Module&) /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:482:13
 #8 0x00007f8d70579a5c HipTextureLoweringPass::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/llvm_passes/HipTextureLowering.cpp:492:10
 #9 0x00007f8d705434a4 llvm::detail::PassModel<llvm::Module, HipTextureLoweringPass, llvm::PreservedAnalyses, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) /gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/include/llvm/IR/PassManagerInternal.h:88:17
#10 0x000000000228812c llvm::PassManager<llvm::Module, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x228812c)
#11 0x000000000078c1a3 llvm::runPassPipeline(llvm::StringRef, llvm::Module&, llvm::TargetMachine*, llvm::TargetLibraryInfoImpl*, llvm::ToolOutputFile*, llvm::ToolOutputFile*, llvm::ToolOutputFile*, llvm::StringRef, llvm::ArrayRef<llvm::StringRef>, llvm::ArrayRef<llvm::PassPlugin>, llvm::opt_tool::OutputKind, llvm::opt_tool::VerifierKind, bool, bool, bool, bool, bool, bool) (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x78c1a3)
#12 0x000000000079cc13 main (/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/opt+0x79cc13)
#13 0x00007f8d708042bd __libc_start_main (/lib64/libc.so.6+0x352bd)
#14 0x0000000000786a2a _start /home/abuild/rpmbuild/BUILD/glibc-2.31/csu/../sysdeps/x86_64/start.S:122:0
clang-15: error: unable to execute command: Aborted (core dumped)
clang-15: error: hipspv-link command failed due to signal (use -v to see invocation)
clang version 15.0.0 (https://github.com/llvm/llvm-project.git 66fa2847a775dda27ddcac3832769441727db42f)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin
clang-15: note: diagnostic msg: Error generating preprocessed source(s).

failed to execute:/gpfs/jlse-fs0/users/pvelesko/install/clang/clang15/clang15-spirv-omp/bin/clang++ -D__HIP_PLATFORM_SPIRV__= -x hip --target=x86_64-linux-gnu -Xclang -no-opaque-pointers --offload=spirv64 -nohipwrapperinc --hip-path=/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build -isystem "/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build/include"  -DKERNELS_PATH="/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/kernels/" -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/external/Catch2 -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/./include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/./kernels -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/build/include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/external/picojson -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/include -I/gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/include/cuspv -Wno-duplicate-decl-specifier -Wno-tautological-constant-compare -Wno-c++20-extensions -Wno-unused-result -Wno-delete-abstract-non-virtual-dtor -Wno-deprecated-declarations -Wunused-command-line-argument -gdwarf-4 --std=c++17 -g -fPIE -Wno-format-extra-args -Wall -pthread -std=c++17 -MD -MT catch/catch_tests/unit/texture/CMakeFiles/TextureTest.dir/hipTextureRef2D.cc.o -MF CMakeFiles/TextureTest.dir/hipTextureRef2D.cc.o.d -o CMakeFiles/TextureTest.dir/hipTextureRef2D.cc.o -c /gpfs/jlse-fs0/users/pvelesko/CHIP-SPV/HIP/tests/catch/unit/texture/hipTextureRef2D.cc
linehill commented 2 years ago

We are dealing with a texture use case that does not lower to SPIR-V straightforwardly. When we lower texture functions to SPIR-V we try to translate their texture object argument to a SPIR-V image type that matches the texture functions’ type (tex1D() -> 1D image, tex2D() -> 2D image, etc.). Here we run into a conflict where the (lowered) texture functions want the texture object to be 1D and 2D image but we can only pick one of them.

We could solve this case by promoting the tex1D() to tex2D() and making the runtime to convert 1D texture object kernel arguments to 2D ones when this kernel is called with some possible loss of performance and additional runtime overhead.

The issue is probably not fixed swiftly and my priorities are elsewhere right now, so I’ll unassign myself from this issue for now.