Fix + improve vprintf implementation (#211)
This commit is contained in:
parent
f0c905db15
commit
fcd7a57888
Binary file not shown.
|
@ -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;
|
||||||
return 2;
|
break;
|
||||||
default:
|
default:
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
return 3;
|
||||||
default:
|
default:
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
return 2;
|
||||||
default:
|
default:
|
||||||
return 0;
|
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)
|
||||||
|
@ -1393,8 +1395,36 @@ extern "C"
|
||||||
char c = *(s++);
|
char c = *(s++);
|
||||||
if (c == 0)
|
if (c == 0)
|
||||||
break;
|
break;
|
||||||
if (c == '%')
|
if (c != '%')
|
||||||
|
continue;
|
||||||
|
|
||||||
|
// %% requires no additional handling
|
||||||
|
if (*s == '%')
|
||||||
{
|
{
|
||||||
|
s++;
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
|
||||||
|
// %s uses __ockl_printf_append_string_n
|
||||||
|
// https://github.com/ROCm/ROCm-Device-Libs/blob/rocm-5.7.x/ockl/src/services.cl#L343
|
||||||
|
if (*s == 's')
|
||||||
|
{
|
||||||
|
s++;
|
||||||
|
const char *value = (const char *)read_valist(valist_ptr, valist_offset, 8);
|
||||||
|
handle = __ockl_printf_append_string_n(handle, value, strlen_plus_one(value), 0);
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
|
||||||
|
// Keep scanning until we figure out the length of this specifier or if we reach the end of the string
|
||||||
|
while (*s != 0) {
|
||||||
|
// "The width is not specified in the format string, but as an additional integer value argument preceding the argument that has to be formatted."
|
||||||
|
if (*s == '*') {
|
||||||
|
s++;
|
||||||
|
uint64_t value = read_valist(valist_ptr, valist_offset, 4);
|
||||||
|
handle = __ockl_printf_append_args(handle, 1, value, 0, 0, 0, 0, 0, 0, 0);
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
|
||||||
uint8_t len = 0;
|
uint8_t len = 0;
|
||||||
if (parse_printf_specifier(s, len))
|
if (parse_printf_specifier(s, len))
|
||||||
{
|
{
|
||||||
|
@ -1406,16 +1436,22 @@ extern "C"
|
||||||
if (specifier_with_length)
|
if (specifier_with_length)
|
||||||
{
|
{
|
||||||
s += specifier_with_length;
|
s += specifier_with_length;
|
||||||
}
|
} else {
|
||||||
if (len > 0)
|
// Assume the unknown character is a sub-specifier and move on
|
||||||
{
|
s++;
|
||||||
uint64_t value = read_valist(valist_ptr, valist_offset, len);
|
continue;
|
||||||
handle = __ockl_printf_append_args(handle, 1, value, 0, 0, 0, 0, 0, 0, 0);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
if (len > 0)
|
||||||
|
{
|
||||||
|
uint64_t value = read_valist(valist_ptr, valist_offset, len);
|
||||||
|
handle = __ockl_printf_append_args(handle, 1, value, 0, 0, 0, 0, 0, 0, 0);
|
||||||
|
}
|
||||||
|
break;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
return (uint32_t)__ockl_printf_append_args(handle, 0, 0, 0, 0, 0, 0, 0, 0, 1);
|
__ockl_printf_append_args(handle, 0, 0, 0, 0, 0, 0, 0, 0, 1);
|
||||||
return 1;
|
return 1;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
Loading…
Reference in a new issue