google / clspv

Clspv is a compiler for OpenCL C to Vulkan compute shaders
Apache License 2.0
623 stars 89 forks source link

[Physical addressing] CTS vec_align_struct triggering assertion #1422

Closed hooneecho closed 3 days ago

hooneecho commented 3 days ago

I came across a compile failure while running one of the OpenCL-CTS suites (./test_vectors vec_align_struct). I was supposed to run it with 64 bit physical address support as below options, but I met an assertion during clBuildProgram.

A kernel which occured the error is like,

typedef struct myUnpackedStruct { 
char c;
    char8 vec;
} testStruct;
__kernel void test_vec_align_struct(__constant char8 *source, __global uint *dest)
{
    __local testStruct test;
    int  tid = get_global_id(0);
    dest[tid] = (uint)((__local uchar *)&(test.vec));
}
While deleting: ptr addrspace(3) %
Use still stuck around after Def is destroyed:  %24 = getelementptr inbounds nuw i8, ptr addrspace(3) <badref>, i64 8
clspv: /home/hooneecho/dev/clspv/third_party/llvm/llvm/lib/IR/Value.cpp:103: llvm::Value::~Value(): Assertion `materialized_use_empty() && "Uses remain when a value is destroyed!"' failed.
hooneecho commented 3 days ago

and reproduce them in compiler explore (left the links below.) (FAIL) with "-arch=spir64 -physical-storage-buffers": https://godbolt.org/z/s4o1jYjPh (SUCCESS) without "-arch=spir64 -physical-storage-buffers": https://godbolt.org/z/a5zseMbbx

when running CLSPV debug version to see the call-stack closely, GV->eraseFromParent(); was seen at lib/LongVectorLoweringPass.cpp:1972 at the very end of the CLSPV. So, I tested again after fixing as below as a workaround, based on just a gut feeling and refering an old issue(https://github.com/google/clspv/pull/1108), and finally the compilation succeeded and the whole test passed.

diff --git a/lib/LongVectorLoweringPass.cpp b/lib/LongVectorLoweringPass.cpp
index 51c30cf4..1a12af07 100644
--- a/lib/LongVectorLoweringPass.cpp
+++ b/lib/LongVectorLoweringPass.cpp
@@ -1969,6 +1969,8 @@ void clspv::LongVectorLoweringPass::cleanDeadGlobals() {
   for (auto const &Mapping : GlobalVariableMap) {
     auto *GV = Mapping.first;
     GV->removeDeadConstantUsers();
-    GV->eraseFromParent();
+    if (GV->getNumUses() == 0) {
+      GV->eraseFromParent();
+    }
   }
 }
hooneecho commented 3 days ago

Working on the above fix by myself