strf::left and strf::right add extra digit in device-side code

[eyalroz]

Consider the following program:

#include <stdio.h>
#include <strf.hpp>

__device__ __host__ void formatting_functions()
{
        char buffer[20];
        auto printer = strf::to(buffer, sizeof(buffer));
        printer ( strf::right(1, 2, '0') );
        printf("%*s\n", 10, buffer);
        printer ( strf::left(1, 2, '0') );
        printf("%*s\n", 10, buffer);
        printer ( strf::right(1, 4, '0') );
        printf("%*s\n", 10, buffer);
        printer ( strf::left(1, 4, '0') );
        printf("%*s\n", 10, buffer);
}

__global__ void ff_kernel()
{
        formatting_functions();
}

int main(void)
{
        int threads_per_block { 1 };
        int blocks_in_grid { 1 };
        printf("Device-side results:\n"); fflush(stdout);
        ff_kernel<<<threads_per_block, blocks_in_grid>>>();
        cudaDeviceReset();
        printf("Host-side results: \n"); fflush(stdout);
        formatting_functions();
}

and let's use the master branch.

When compiling this, we get warnings about calling a host-side from device-side functions; that's the first problem. But - it does build. And the output is:

Device-side results:
       001
       100
     00001
     10000
Host-side results: 
        01
        10
      0001
      1000

So, an extra fill digits. That's the bug.

[eyalroz]

Also have a look at this commit on my repository:

• It exposes this bug.
• If you uncomment some more code there, you get an illegal access error, which I am 90% certain is within strf code, not due to the text itself.

[robhz786]

I will take a look at that.
Meanwhile, if you need to fill with '0', you can use the 'p(precision)' format function instead. This one seems to be working.

[robhz786]

The bug is fixed on branch hotfix-0.10.4. It has to do with the max function on CUDA_ARCH.
I still don't know what causes the other bug ( the illegal access error ). But there is a workaround:

__global__ void formatting_functions(char* buffer, std::size_t buffer_size)
{
    strf::cstr_writer writer(buffer, buffer_size);
    auto printer = strf::to(writer);

    printer ("strf::fmt(0) gives ", strf::fmt(0), '\n');
    printer ("strf::fmt(0).hex() gives ", strf::fmt(0).hex(), '\n');
    printer ("strf::fmt(0).bin() gives ", strf::fmt(0).bin(), '\n');
    printer ("strf::left(0, 2, '0') gives ", strf::left(0, 2, '0'), '\n');
    printer ("strf::right(0, 2, '0') gives ", strf::right(0, 2, '0'), '\n');
    printer ("strf::fmt(123) gives ", strf::fmt(123), '\n');
    printer ("strf::fmt(123).hex() gives ", strf::fmt(123).hex(), '\n');
    printer ("strf::fmt(123).bin() gives ", strf::fmt(123).bin(), '\n');
    printer ("strf::left(123, 5, '0') gives ", strf::left(123, 5, '0'), '\n');
    printer ("strf::right(123, 5, '0') gives ", strf::right(123, 5, '0'), '\n');

    writer.finish();
}

I commited that as well.

[eyalroz]

Thanks. I made a few comments on the 0.10.4 commits.