Fix + improve vprintf implementation

[NyanCatTW1]

We were trying to debug a kernel with printf, but then we realized %d and friends didn't work at all! So we dug in and made changes to vprintf that should allow it to support all sorts of common printf calls

Test code (mostly taken from https://cplusplus.com/reference/cstdio/printf/)

#include <stdio.h>
#include <cuda.h>

__global__ void kernel() {
    printf("Characters: %c %c \n", 'a', 65);
    printf("Decimals: %d %ld\n", 1977, 650000L);
    printf("Preceding with blanks: %10d \n", 1977);
    printf("Preceding with zeros: %010d \n", 1977);
    printf("Some different radices: %d %x %o %#x %#o \n", 100, 100, 100, 100, 100);
    printf("floats: %4.2f %+.0e %E \n", 3.1416, 3.1416, 3.1416);
    printf("Width trick: %*d \n", 5, 10);
    printf("%s \n", "A string");
    printf("I'm not formatting anything: %%s!\n");
}

Expected output (currently identical to CUDA's)

Characters: a A 
Decimals: 1977 650000
Preceding with blanks:       1977 
Preceding with zeros: 0000001977 
Some different radices: 100 64 144 0x64 0144 
floats: 3.14 +3e+00 3.141600E+00 
Width trick:    10 
A string 
I'm not formatting anything: %s!

Some of the code could've been done better, but well, at least now it works properly

[vosen]

Mostly good, but I need some small changes

[vosen]

Some of the code could've been done better, but well, at least now it works properly

Yes, te vprintf code is not very good. And difficult to test properly: last time I tried to capture from stdout handle it didn't work with HIP.
I should probably split printf into a separate file and test it on CPU with mocked __ockl_printf_*, but it's a story for another day.