mirrors/ZLUDA
0
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-08-05 15:49:24 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

Fix vprintf implementation

Browse source
This commit is contained in:
NyanCatTW1 2024-04-23 20:43:26 +03:00
commit f7534750c2
2 changed files with 11 additions and 9 deletions
Download patch file Download diff file Expand all files Collapse all files

BIN
ptx/lib/zluda_ptx_impl.bc
View file

Binary file not shown.

22
ptx/lib/zluda_ptx_impl.cpp
View file

@ -1,5 +1,5 @@
// Compile and disassemble: // Compile and disassemble:
// python3 ./cvt.py > cvt.h && /opt/rocm/llvm/bin/clang -std=c++17 -Xclang -no-opaque-pointers -Wall -Wextra -Wsign-compare -Wconversion -x hip zluda_ptx_impl.cpp -S -emit-llvm --cuda-device-only -nogpulib -O3 -Xclang -fallow-half-arguments-and-returns -o - | sed -e 's/define/define linkonce_odr/g' | sed -e '/@llvm.used/d' | sed -e 's/\"target-cpu\"=\"[^\"]*\"//g' | sed -e 's/\"target-features\"=\"[^\"]*\"//g' | sed -e 's/\"denormal-fp-math-f32\"=\"[^\"]*\"//g' | sed -e 's/!llvm.module.flags = !{!0, !1, !2, !3, !4}/!llvm.module.flags = !{ }/g' | sed -e 's/memory(none)/readnone/g' | sed -e 's/memory(argmem: readwrite, inaccessiblemem: readwrite)/inaccessiblemem_or_argmemonly/g' | sed -e 's/memory(read)/readonly/g' | sed -e 's/memory(argmem: readwrite)/argmemonly/g' | llvm-as-13 -o zluda_ptx_impl.bc && /opt/rocm/llvm/bin/llvm-dis zluda_ptx_impl.bc // python3 ./cvt.py > cvt.h && /opt/rocm/llvm/bin/clang -std=c++20 -Xclang -no-opaque-pointers -Wall -Wextra -Wsign-compare -Wconversion -x hip zluda_ptx_impl.cpp -S -emit-llvm --cuda-device-only -nogpulib -O3 -Xclang -fallow-half-arguments-and-returns -o - | sed -e 's/define/define linkonce_odr/g' | sed -e '/@llvm.used/d' | sed -e 's/\"target-cpu\"=\"[^\"]*\"//g' | sed -e 's/\"target-features\"=\"[^\"]*\"//g' | sed -e 's/\"denormal-fp-math-f32\"=\"[^\"]*\"//g' | sed -e 's/!llvm.module.flags = !{!0, !1, !2, !3, !4}/!llvm.module.flags = !{ }/g' | sed -e 's/memory(none)/readnone/g' | sed -e 's/memory(argmem: readwrite, inaccessiblemem: readwrite)/inaccessiblemem_or_argmemonly/g' | sed -e 's/memory(read)/readonly/g' | sed -e 's/memory(argmem: readwrite)/argmemonly/g' | llvm-as-13 -o zluda_ptx_impl.bc && /opt/rocm/llvm/bin/llvm-dis zluda_ptx_impl.bc
// Compile to binary: // Compile to binary:
// /opt/rocm/llvm/bin/clang -x ir -target amdgcn-amd-amdhsa -Xlinker --no-undefined zluda_ptx_impl.bc -mno-wavefrontsize64 -mcpu=gfx1030 // /opt/rocm/llvm/bin/clang -x ir -target amdgcn-amd-amdhsa -Xlinker --no-undefined zluda_ptx_impl.bc -mno-wavefrontsize64 -mcpu=gfx1030
// Decompile: // Decompile:
@ -1260,6 +1260,7 @@ extern "C"
default: default:
return 0; return 0;
} }
return 2;
case 'l': case 'l':
switch (s[1]) switch (s[1])
{ {
@ -1289,17 +1290,18 @@ extern "C"
case 'X': case 'X':
case 'n': case 'n':
len = 8; len = 8;
break;
default:
return 0;
}
return 3;
default:
return 0;
}
return 2; return 2;
default: default:
return 0; return 0;
} }
default:
return 0;
}
default:
return 0;
}
return 1;
} }
__device__ static bool parse_printf_specifier(const char *s, uint8_t &len) __device__ static bool parse_printf_specifier(const char *s, uint8_t &len)
@ -1407,6 +1409,8 @@ extern "C"
{ {
s += specifier_with_length; s += specifier_with_length;
} }
}
if (len > 0) if (len > 0)
{ {
uint64_t value = read_valist(valist_ptr, valist_offset, len); uint64_t value = read_valist(valist_ptr, valist_offset, len);
@ -1414,9 +1418,7 @@ extern "C"
} }
} }
} }
}
return (uint32_t)__ockl_printf_append_args(handle, 0, 0, 0, 0, 0, 0, 0, 0, 1); return (uint32_t)__ockl_printf_append_args(handle, 0, 0, 0, 0, 0, 0, 0, 0, 1);
return 1;
} }
int64_t __ockl_mul_hi_i64(int64_t x, int64_t y) __attribute__((device)); int64_t __ockl_mul_hi_i64(int64_t x, int64_t y) __attribute__((device));