diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 62d69fc3..660f6163 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1,24 +1,63 @@ -cmake_minimum_required(VERSION 3.9) +cmake_minimum_required(VERSION 3.18) enable_language(CXX) -add_executable(test_suite "test_suite.cpp") - -set_target_properties( - test_suite - PROPERTIES - CXX_STANDARD 11 - CXX_STANDARD_REQUIRED YES - CXX_EXTENSIONS NO -) - +option(BUILD_CUDA_TESTS "Include tests for use of the printf library with CUDA host-side and device-side code" OFF) option(TEST_BROKEN_FORMATS "Include tests using non-standard-compliant format strings?" ON) -# ... don't worry, we'll suppress the compiler warnings for those. -if (TEST_BROKEN_FORMATS) - target_compile_definitions(test_suite PRIVATE TEST_WITH_NON_STANDARD_FORMAT_STRINGS) + +set(cuda_tests) +if(BUILD_CUDA_TESTS) + enable_language(CUDA) + list(APPEND CMAKE_CUDA_FLAGS "--extended-lambda --expt-relaxed-constexpr -Xcudafe --display_error_number") + if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.18) + cmake_policy(SET CMP0104 OLD) + endif() + include(FindCUDA/select_compute_arch) + list(APPEND cuda_tests cuda_test_suite_host cuda_test_suite_device) + if((NOT DEFINED CUDA_ARCH_FLAGS) OR ("${CUDA_ARCH_FLAGS}" STREQUAL "")) + cuda_select_nvcc_arch_flags(CUDA_ARCH_FLAGS_1 Auto) + set(CUDA_ARCH_FLAGS ${CUDA_ARCH_FLAGS} CACHE STRING "CUDA -gencode parameters") + string(REPLACE ";" " " CUDA_ARCH_FLAGS_STR "${CUDA_ARCH_FLAGS}") + else() + set(CUDA_ARCH_FLAGS_STR "${CUDA_ARCH_FLAGS}") + endif() + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} ${CUDA_ARCH_FLAGS_STR}") endif() -target_link_libraries(test_suite PRIVATE mpaland-printf) +foreach(tgt ${cuda_tests}) + add_executable(${tgt} "${tgt}.cu") + set_target_properties( + ${tgt} + PROPERTIES + CUDA_STANDARD 11 + CUDA_STANDARD_REQUIRED YES + CUDA_EXTENSIONS NO + ) +endforeach() + +set(tests test_suite ${cuda_tests}) + +add_executable(test_suite "test_suite.cpp") + +foreach(tgt ${tests}) + set_target_properties( + ${tgt} + PROPERTIES + CXX_STANDARD 11 + CXX_STANDARD_REQUIRED YES + CXX_EXTENSIONS NO + ) + if (TEST_BROKEN_FORMATS) + # ... don't worry, we'll suppress the compiler warnings for those. + target_compile_definitions(${tgt} PRIVATE TEST_WITH_NON_STANDARD_FORMAT_STRINGS) + endif() + + target_link_libraries(test_suite PRIVATE mpaland-printf) + add_test( + NAME ${PROJECT_NAME}.${tgt} + COMMAND ${tgt} # ${TEST_RUNNER_PARAMS} + ) +endforeach() if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC") target_compile_options(test_suite PRIVATE /W4) @@ -53,9 +92,5 @@ elseif (CMAKE_CXX_COMPILER_ID STREQUAL "GNU" OR endif() endif() -add_test( - NAME ${PROJECT_NAME}.test_suite - COMMAND test_suite # ${TEST_RUNNER_PARAMS} -) diff --git a/test/cuda_test_suite_device.cu b/test/cuda_test_suite_device.cu new file mode 100644 index 00000000..9c8e4e99 --- /dev/null +++ b/test/cuda_test_suite_device.cu @@ -0,0 +1,1798 @@ +#define CATCH_CONFIG_MAIN +#include "catch.hpp" + +#include +#include +#include +#include +#include + +namespace test { + // use functions in own test namespace to avoid stdio conflicts + #include "../printf.h" + #include "../printf.c" +} // namespace test + +#undef printf + +// Multi-compiler-compatible local warning suppression + +#if defined(_MSC_VER) + #define DISABLE_WARNING_PUSH __pragma(warning( push )) + #define DISABLE_WARNING_POP __pragma(warning( pop )) + #define DISABLE_WARNING(warningNumber) __pragma(warning( disable : warningNumber )) + + // TODO: find the right warning number for this + #define DISABLE_WARNING_PRINTF_FORMAT + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + +#elif defined(__NVCC__) + #define DO_PRAGMA(X) _Pragma(#X) + #define DISABLE_WARNING_PUSH DO_PRAGMA(push) + #define DISABLE_WARNING_POP DO_PRAGMA(pop) + #define DISABLE_WARNING(warning_code) DO_PRAGMA(diag_suppress warning_code) + + #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(bad_printf_format_string) + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + +#elif defined(__GNUC__) || defined(__clang__) + #define DO_PRAGMA(X) _Pragma(#X) + #define DISABLE_WARNING_PUSH DO_PRAGMA(GCC diagnostic push) + #define DISABLE_WARNING_POP DO_PRAGMA(GCC diagnostic pop) + #define DISABLE_WARNING(warningName) DO_PRAGMA(GCC diagnostic ignored #warningName) + + #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(-Wformat) + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS DISABLE_WARNING(-Wformat-extra-args) + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW DISABLE_WARNING(-Wformat-overflow) +#else + #define DISABLE_WARNING_PUSH + #define DISABLE_WARNING_POP + #define DISABLE_WARNING_PRINTF_FORMAT + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW +#endif + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +DISABLE_WARNING_PUSH +DISABLE_WARNING_PRINTF_FORMAT +DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS +#endif + +char* make_device_string(char const* s) +{ + size_t size = strlen(s) + 1; + void* dsptr; + cudaMalloc(&dsptr, size); + cudaMemcpy(dsptr, s, size, cudaMemcpyDefault); + cudaDeviceSynchronize(); + return (char *) dsptr; +} + +inline char* mds(char const* s) { return make_device_string(s); } + +struct poor_mans_string_view { + char* data; + size_t size; +}; + +struct sv_and_pos { + const poor_mans_string_view sv; + size_t pos; +}; + +__device__ void append_to_buffer(char c, void* type_erased_svnp) +{ + auto& svnp = *(static_cast(type_erased_svnp)); + if (svnp.pos < svnp.sv.size) { + svnp.sv.data[svnp.pos++] = c; + } +} + +// output function type +typedef void (*out_fct_type)(char character, void* arg); + +// ... just need to make the linker happy :-( +PRINTF_HOST void test::_putchar(char character) +{ + exit(EXIT_FAILURE); +} + +enum class invokable { + sprintf, vsprintf, snprintf, vsnprintf +}; + +__device__ __host__ char const* name(invokable inv) +{ + switch(inv) { + case invokable::sprintf: return "sprintf"; + case invokable::snprintf: return "snprintf"; + case invokable::vsprintf: return "vsprintf"; + case invokable::vsnprintf: return "vsnprintf"; + } + return "unknown"; +} + +__device__ int vsprintf_wrapper(char* buffer, char const* format, ...) +{ + va_list args; + va_start(args, format); + int ret = test::vsprintf(buffer, format, args); + va_end(args); + return ret; +} + +__device__ int vnsprintf_wrapper(char* buffer, size_t buffer_size, char const* format, ...) +{ + va_list args; + va_start(args, format); + int ret = test::vsnprintf(buffer, buffer_size, format, args); + va_end(args); + return ret; +} + +namespace kernels { + +template +__global__ void +invoke( + int * __restrict__ result, + invokable which, + char * __restrict__ buffer, + size_t buffer_size, + char const * __restrict__ format, + Ts... args) +{ + switch(which) { + case invokable::sprintf: *result = test::sprintf(buffer, format, args...); break; + case invokable::snprintf: *result = test::snprintf(buffer, buffer_size, format, args...); break; + case invokable::vsprintf: *result = vsprintf_wrapper(buffer, format, args...); break; + case invokable::vsnprintf: *result = vnsprintf_wrapper(buffer, buffer_size, format, args...); break; + } +} + +} // namespace kernels + +template +int invoke_on_device(invokable which, char* buffer, size_t buffer_size, char const* format, Ts... args) +{ + char* buffer_d; + char* format_d; + int* result_d; + int result; + size_t format_size = strlen(format) + 1; + cudaGetLastError(); // Clearing/ignoring earlier errors + cudaMalloc(&result_d, sizeof(int)); + if (buffer != nullptr) { + cudaMalloc(&buffer_d, buffer_size); + cudaMemcpy(buffer_d, buffer, buffer_size, cudaMemcpyDefault); + } else { + buffer_d = nullptr ; + } + if (format != nullptr) { + cudaMalloc(&format_d, format_size); + cudaMemcpy(format_d, format, format_size, cudaMemcpyDefault); + } + else { + format_d = nullptr; + } + kernels::invoke<<<1, 1>>>(result_d, which, buffer_d, buffer_size, format_d, args...); // Note: No perfect forwarding. + cudaDeviceSynchronize(); + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + throw std::runtime_error("CUDA error: " + std::string(cudaGetErrorString(error))); + } + if (buffer != nullptr) { + cudaMemcpy(buffer, buffer_d, buffer_size, cudaMemcpyDefault); + } + cudaMemcpy(&result, result_d, sizeof(int), cudaMemcpyDefault); + cudaFree(buffer_d); + cudaFree(format_d); + cudaFree(result_d); + cudaDeviceSynchronize(); + error = cudaGetLastError(); + if (error != cudaSuccess) { + throw std::runtime_error("CUDA error: " + std::string(cudaGetErrorString(error))); + } + return result; +} + +constexpr const size_t base_buffer_size { 100 }; + +template +int invoke_on_device(invokable which, char* buffer, char const* format, Ts... args) +{ + return invoke_on_device(which, buffer, base_buffer_size, format, args...); +} + +namespace kernels { + +__global__ void fctprintf(char* buffer) +{ + sv_and_pos svnp { {buffer, base_buffer_size}, 0 }; + test::fctprintf(append_to_buffer, &svnp, "This is a test of %X", 0x12EFU); +} + +} // namespace kernels + +TEST_CASE("fctprintf", "[]" ) { + char buffer[base_buffer_size]; + char* buffer_d; + cudaMalloc(&buffer_d, base_buffer_size); + cudaMemset(buffer_d, 0xCC, base_buffer_size); + kernels::fctprintf<<<1, 1>>>(buffer_d); + cudaMemcpy(buffer, buffer_d, base_buffer_size, cudaMemcpyDefault); + cudaDeviceSynchronize(); + REQUIRE(!strncmp(buffer, "This is a test of 12EF", 22U)); + // Remember: printf does not append a `\0` to the output after going through its format string. + REQUIRE(buffer[22] == (char)0xCC); + cudaFree(buffer_d); +} + +PRINTF_HD static void vfctprintf_builder_1(out_fct_type f, void* f_arg, ...) +{ + va_list args; + va_start(args, f_arg); + test::vfctprintf(f, f_arg, "This is a test of %X", args); + va_end(args); +} + +namespace kernels { + +__global__ void vfctprintf(char* buffer) +{ + sv_and_pos svnp { {buffer, base_buffer_size}, 0 }; + vfctprintf_builder_1(append_to_buffer, &svnp, 0x12EFU); +} + +} // namespace kernels + +TEST_CASE("vfctprintf", "[]" ) { + char buffer[base_buffer_size]; + char* buffer_d; + cudaMalloc(&buffer_d, base_buffer_size); + cudaMemset(buffer_d, 0xCC, base_buffer_size); + kernels::vfctprintf<<<1, 1>>>(buffer_d); + cudaMemcpy(buffer, buffer_d, base_buffer_size, cudaMemcpyDefault); + cudaDeviceSynchronize(); + REQUIRE(!strncmp(buffer, "This is a test of 12EF", 22U)); + REQUIRE(buffer[22] == (char)0xCC); + cudaFree(buffer_d); +} + +namespace kernels { + +__global__ void snprintf(char* buffer, size_t buffer_size) +{ + test::snprintf(buffer, buffer_size, "%d", -1000); +} + +} // namespace kernels + +TEST_CASE("snprintf", "[]" ) { + char buffer[base_buffer_size]; + invoke_on_device(invokable::snprintf, buffer, "%d", -1000); + REQUIRE(!strcmp(buffer, "-1000")); + invoke_on_device(invokable::snprintf, buffer, 3, "%d", -1000); + REQUIRE(!strcmp(buffer, "-1")); +} + +TEST_CASE("vsprintf", "[]" ) { + char buffer[base_buffer_size]; + invoke_on_device(invokable::vsprintf, buffer, "%d", -1 ); + REQUIRE(!strcmp(buffer, "-1")); + + invoke_on_device(invokable::vsprintf, buffer, "%d %d %s", 3, -1000, mds("test") ); + REQUIRE(!strcmp(buffer, "3 -1000 test")); +} + +TEST_CASE("vsnprintf", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::vsnprintf, buffer, "%d", -1); + REQUIRE(!strcmp(buffer, "-1")); + + invoke_on_device(invokable::vsnprintf, buffer, "%d %d %s", 3, -1000, mds("test")); + REQUIRE(!strcmp(buffer, "3 -1000 test")); +} + +TEST_CASE("simple sprintf", "[]" ) { + char buffer[base_buffer_size]; + memset(buffer, 0xCC, base_buffer_size); + invoke_on_device(invokable::sprintf, buffer, "%d", 42); + REQUIRE(!strcmp(buffer, "42")); +} + +TEST_CASE("space flag", "[]" ) { + char buffer[base_buffer_size]; + memset(buffer, 0xCC, base_buffer_size); + + invoke_on_device(invokable::sprintf, buffer, "% d", 42); + REQUIRE(!strcmp(buffer, " 42")); + + invoke_on_device(invokable::sprintf, buffer, "% d", -42); + REQUIRE(!strcmp(buffer, "-42")); + + invoke_on_device(invokable::sprintf, buffer, "% 5d", 42); + REQUIRE(!strcmp(buffer, " 42")); + + invoke_on_device(invokable::sprintf, buffer, "% 5d", -42); + REQUIRE(!strcmp(buffer, " -42")); + + invoke_on_device(invokable::sprintf, buffer, "% 15d", 42); + REQUIRE(!strcmp(buffer, " 42")); + + invoke_on_device(invokable::sprintf, buffer, "% 15d", -42); + REQUIRE(!strcmp(buffer, " -42")); + + invoke_on_device(invokable::sprintf, buffer, "% 15d", -42); + REQUIRE(!strcmp(buffer, " -42")); + + invoke_on_device(invokable::sprintf, buffer, "% 15.3f", -42.987); + REQUIRE(!strcmp(buffer, " -42.987")); + + invoke_on_device(invokable::sprintf, buffer, "% 15.3f", 42.987); + REQUIRE(!strcmp(buffer, " 42.987")); + + invoke_on_device(invokable::sprintf, buffer, "% d", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "% d", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "% i", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "% i", -1024); + REQUIRE(!strcmp(buffer, "-1024")); +} + + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("space flag - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "% s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, "Hello testing")); + + invoke_on_device(invokable::sprintf, buffer, "% u", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "% u", 4294966272U); + REQUIRE(!strcmp(buffer, "4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "% o", 511); + REQUIRE(!strcmp(buffer, "777")); + + invoke_on_device(invokable::sprintf, buffer, "% o", 4294966785U); + REQUIRE(!strcmp(buffer, "37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "% x", 305441741); + REQUIRE(!strcmp(buffer, "1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "% x", 3989525555U); + REQUIRE(!strcmp(buffer, "edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "% X", 305441741); + REQUIRE(!strcmp(buffer, "1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "% X", 3989525555U); + REQUIRE(!strcmp(buffer, "EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "% c", 'x'); + REQUIRE(!strcmp(buffer, "x")); +} +#endif + +TEST_CASE("+ flag", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%+d", 42); + REQUIRE(!strcmp(buffer, "+42")); + + invoke_on_device(invokable::sprintf, buffer, "%+d", -42); + REQUIRE(!strcmp(buffer, "-42")); + + invoke_on_device(invokable::sprintf, buffer, "%+5d", 42); + REQUIRE(!strcmp(buffer, " +42")); + + invoke_on_device(invokable::sprintf, buffer, "%+5d", -42); + REQUIRE(!strcmp(buffer, " -42")); + + invoke_on_device(invokable::sprintf, buffer, "%+15d", 42); + REQUIRE(!strcmp(buffer, " +42")); + + invoke_on_device(invokable::sprintf, buffer, "%+15d", -42); + REQUIRE(!strcmp(buffer, " -42")); + + invoke_on_device(invokable::sprintf, buffer, "%+d", 1024); + REQUIRE(!strcmp(buffer, "+1024")); + + invoke_on_device(invokable::sprintf, buffer, "%+d", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "%+i", 1024); + REQUIRE(!strcmp(buffer, "+1024")); + + invoke_on_device(invokable::sprintf, buffer, "%+i", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "%+.0d", 0); + REQUIRE(!strcmp(buffer, "+")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("+ flag - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%+s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, "Hello testing")); + + invoke_on_device(invokable::sprintf, buffer, "%+u", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%+u", 4294966272U); + REQUIRE(!strcmp(buffer, "4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%+o", 511); + REQUIRE(!strcmp(buffer, "777")); + + invoke_on_device(invokable::sprintf, buffer, "%+o", 4294966785U); + REQUIRE(!strcmp(buffer, "37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%+x", 305441741); + REQUIRE(!strcmp(buffer, "1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%+x", 3989525555U); + REQUIRE(!strcmp(buffer, "edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%+X", 305441741); + REQUIRE(!strcmp(buffer, "1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%+X", 3989525555U); + REQUIRE(!strcmp(buffer, "EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "%+c", 'x'); + REQUIRE(!strcmp(buffer, "x")); +} +#endif + + +TEST_CASE("0 flag", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%0d", 42); + REQUIRE(!strcmp(buffer, "42")); + + invoke_on_device(invokable::sprintf, buffer, "%0ld", 42L); + REQUIRE(!strcmp(buffer, "42")); + + invoke_on_device(invokable::sprintf, buffer, "%0d", -42); + REQUIRE(!strcmp(buffer, "-42")); + + invoke_on_device(invokable::sprintf, buffer, "%05d", 42); + REQUIRE(!strcmp(buffer, "00042")); + + invoke_on_device(invokable::sprintf, buffer, "%05d", -42); + REQUIRE(!strcmp(buffer, "-0042")); + + invoke_on_device(invokable::sprintf, buffer, "%015d", 42); + REQUIRE(!strcmp(buffer, "000000000000042")); + + invoke_on_device(invokable::sprintf, buffer, "%015d", -42); + REQUIRE(!strcmp(buffer, "-00000000000042")); + + invoke_on_device(invokable::sprintf, buffer, "%015.2f", 42.1234); + REQUIRE(!strcmp(buffer, "000000000042.12")); + + invoke_on_device(invokable::sprintf, buffer, "%015.3f", 42.9876); + REQUIRE(!strcmp(buffer, "00000000042.988")); + + invoke_on_device(invokable::sprintf, buffer, "%015.5f", -42.9876); + REQUIRE(!strcmp(buffer, "-00000042.98760")); +} + + +TEST_CASE("- flag", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%-d", 42); + REQUIRE(!strcmp(buffer, "42")); + + invoke_on_device(invokable::sprintf, buffer, "%-d", -42); + REQUIRE(!strcmp(buffer, "-42")); + + invoke_on_device(invokable::sprintf, buffer, "%-5d", 42); + REQUIRE(!strcmp(buffer, "42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-5d", -42); + REQUIRE(!strcmp(buffer, "-42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-15d", 42); + REQUIRE(!strcmp(buffer, "42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-15d", -42); + REQUIRE(!strcmp(buffer, "-42 ")); +} + + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("- flag - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%-0d", 42); + REQUIRE(!strcmp(buffer, "42")); + + invoke_on_device(invokable::sprintf, buffer, "%-0d", -42); + REQUIRE(!strcmp(buffer, "-42")); + + invoke_on_device(invokable::sprintf, buffer, "%-05d", 42); + REQUIRE(!strcmp(buffer, "42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-05d", -42); + REQUIRE(!strcmp(buffer, "-42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-015d", 42); + REQUIRE(!strcmp(buffer, "42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-015d", -42); + REQUIRE(!strcmp(buffer, "-42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-d", 42); + REQUIRE(!strcmp(buffer, "42")); + + invoke_on_device(invokable::sprintf, buffer, "%0-d", -42); + REQUIRE(!strcmp(buffer, "-42")); + + invoke_on_device(invokable::sprintf, buffer, "%0-5d", 42); + REQUIRE(!strcmp(buffer, "42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-5d", -42); + REQUIRE(!strcmp(buffer, "-42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-15d", 42); + REQUIRE(!strcmp(buffer, "42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-15d", -42); + REQUIRE(!strcmp(buffer, "-42 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-15.3e", -42.); +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + REQUIRE(!strcmp(buffer, "-4.200e+01 ")); +#else + REQUIRE(!strcmp(buffer, "e")); +#endif + + invoke_on_device(invokable::sprintf, buffer, "%0-15.3g", -42.); +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + REQUIRE(!strcmp(buffer, "-42.0 ")); +#else + REQUIRE(!strcmp(buffer, "g")); +#endif +} +#endif + + +TEST_CASE("# flag", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%#.0x", 0); + REQUIRE(!strcmp(buffer, "")); + invoke_on_device(invokable::sprintf, buffer, "%#.1x", 0); + REQUIRE(!strcmp(buffer, "0")); + invoke_on_device(invokable::sprintf, buffer, "%#.0llx", (long long)0); + REQUIRE(!strcmp(buffer, "")); + invoke_on_device(invokable::sprintf, buffer, "%#.8x", 0x614e); + REQUIRE(!strcmp(buffer, "0x0000614e")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("# flag - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer,"%#b", 6); + REQUIRE(!strcmp(buffer, "0b110")); +} +#endif + +TEST_CASE("specifier", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "Hello testing"); + REQUIRE(!strcmp(buffer, "Hello testing")); + + invoke_on_device(invokable::sprintf, buffer, "%s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, "Hello testing")); + +DISABLE_WARNING_PUSH +DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + invoke_on_device(invokable::sprintf, buffer, "%s", NULL); +DISABLE_WARNING_POP + REQUIRE(!strcmp(buffer, "(null)")); + + invoke_on_device(invokable::sprintf, buffer, "%d", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%d", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "%i", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%i", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "%u", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%u", 4294966272U); + REQUIRE(!strcmp(buffer, "4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%o", 511); + REQUIRE(!strcmp(buffer, "777")); + + invoke_on_device(invokable::sprintf, buffer, "%o", 4294966785U); + REQUIRE(!strcmp(buffer, "37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%x", 305441741); + REQUIRE(!strcmp(buffer, "1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%x", 3989525555U); + REQUIRE(!strcmp(buffer, "edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%X", 305441741); + REQUIRE(!strcmp(buffer, "1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%X", 3989525555U); + REQUIRE(!strcmp(buffer, "EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "%%"); + REQUIRE(!strcmp(buffer, "%")); +} + + +TEST_CASE("width", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%1s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, "Hello testing")); + + invoke_on_device(invokable::sprintf, buffer, "%1d", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%1d", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "%1i", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%1i", -1024); + REQUIRE(!strcmp(buffer, "-1024")); + + invoke_on_device(invokable::sprintf, buffer, "%1u", 1024); + REQUIRE(!strcmp(buffer, "1024")); + + invoke_on_device(invokable::sprintf, buffer, "%1u", 4294966272U); + REQUIRE(!strcmp(buffer, "4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%1o", 511); + REQUIRE(!strcmp(buffer, "777")); + + invoke_on_device(invokable::sprintf, buffer, "%1o", 4294966785U); + REQUIRE(!strcmp(buffer, "37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%1x", 305441741); + REQUIRE(!strcmp(buffer, "1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%1x", 3989525555U); + REQUIRE(!strcmp(buffer, "edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%1X", 305441741); + REQUIRE(!strcmp(buffer, "1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%1X", 3989525555U); + REQUIRE(!strcmp(buffer, "EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "%1c", 'x'); + REQUIRE(!strcmp(buffer, "x")); +} + + +TEST_CASE("width 20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%20s", mds("Hello")); + REQUIRE(!strcmp(buffer, " Hello")); + + invoke_on_device(invokable::sprintf, buffer, "%20d", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20d", -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20i", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20i", -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20u", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20u", 4294966272U); + REQUIRE(!strcmp(buffer, " 4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%20o", 511); + REQUIRE(!strcmp(buffer, " 777")); + + invoke_on_device(invokable::sprintf, buffer, "%20o", 4294966785U); + REQUIRE(!strcmp(buffer, " 37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%20x", 305441741); + REQUIRE(!strcmp(buffer, " 1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%20x", 3989525555U); + REQUIRE(!strcmp(buffer, " edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%20X", 305441741); + REQUIRE(!strcmp(buffer, " 1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%20X", 3989525555U); + REQUIRE(!strcmp(buffer, " EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "%20c", 'x'); + REQUIRE(!strcmp(buffer, " x")); +} + + +TEST_CASE("width *20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%*s", 20, mds("Hello")); + REQUIRE(!strcmp(buffer, " Hello")); + + invoke_on_device(invokable::sprintf, buffer, "%*d", 20, 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%*d", 20, -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%*i", 20, 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%*i", 20, -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%*u", 20, 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%*u", 20, 4294966272U); + REQUIRE(!strcmp(buffer, " 4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%*o", 20, 511); + REQUIRE(!strcmp(buffer, " 777")); + + invoke_on_device(invokable::sprintf, buffer, "%*o", 20, 4294966785U); + REQUIRE(!strcmp(buffer, " 37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%*x", 20, 305441741); + REQUIRE(!strcmp(buffer, " 1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%*x", 20, 3989525555U); + REQUIRE(!strcmp(buffer, " edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%*X", 20, 305441741); + REQUIRE(!strcmp(buffer, " 1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%*X", 20, 3989525555U); + REQUIRE(!strcmp(buffer, " EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "%*c", 20,'x'); + REQUIRE(!strcmp(buffer, " x")); +} + + +TEST_CASE("width -20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%-20s", mds("Hello")); + REQUIRE(!strcmp(buffer, "Hello ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20d", 1024); + REQUIRE(!strcmp(buffer, "1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20d", -1024); + REQUIRE(!strcmp(buffer, "-1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20i", 1024); + REQUIRE(!strcmp(buffer, "1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20i", -1024); + REQUIRE(!strcmp(buffer, "-1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20u", 1024); + REQUIRE(!strcmp(buffer, "1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20.4f", 1024.1234); + REQUIRE(!strcmp(buffer, "1024.1234 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20u", 4294966272U); + REQUIRE(!strcmp(buffer, "4294966272 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20o", 511); + REQUIRE(!strcmp(buffer, "777 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20o", 4294966785U); + REQUIRE(!strcmp(buffer, "37777777001 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20x", 305441741); + REQUIRE(!strcmp(buffer, "1234abcd ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20x", 3989525555U); + REQUIRE(!strcmp(buffer, "edcb5433 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20X", 305441741); + REQUIRE(!strcmp(buffer, "1234ABCD ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20X", 3989525555U); + REQUIRE(!strcmp(buffer, "EDCB5433 ")); + + invoke_on_device(invokable::sprintf, buffer, "%-20c", 'x'); + REQUIRE(!strcmp(buffer, "x ")); + + invoke_on_device(invokable::sprintf, buffer, "|%5d| |%-2d| |%5d|", 9, 9, 9); + REQUIRE(!strcmp(buffer, "| 9| |9 | | 9|")); + + invoke_on_device(invokable::sprintf, buffer, "|%5d| |%-2d| |%5d|", 10, 10, 10); + REQUIRE(!strcmp(buffer, "| 10| |10| | 10|")); + + invoke_on_device(invokable::sprintf, buffer, "|%5d| |%-12d| |%5d|", 9, 9, 9); + REQUIRE(!strcmp(buffer, "| 9| |9 | | 9|")); + + invoke_on_device(invokable::sprintf, buffer, "|%5d| |%-12d| |%5d|", 10, 10, 10); + REQUIRE(!strcmp(buffer, "| 10| |10 | | 10|")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("width 0-20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%0-20s", mds("Hello")); + REQUIRE(!strcmp(buffer, "Hello ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20d", 1024); + REQUIRE(!strcmp(buffer, "1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20d", -1024); + REQUIRE(!strcmp(buffer, "-1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20i", 1024); + REQUIRE(!strcmp(buffer, "1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20i", -1024); + REQUIRE(!strcmp(buffer, "-1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20u", 1024); + REQUIRE(!strcmp(buffer, "1024 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20u", 4294966272U); + REQUIRE(!strcmp(buffer, "4294966272 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20o", 511); + REQUIRE(!strcmp(buffer, "777 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20o", 4294966785U); + REQUIRE(!strcmp(buffer, "37777777001 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20x", 305441741); + REQUIRE(!strcmp(buffer, "1234abcd ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20x", 3989525555U); + REQUIRE(!strcmp(buffer, "edcb5433 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20X", 305441741); + REQUIRE(!strcmp(buffer, "1234ABCD ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20X", 3989525555U); + REQUIRE(!strcmp(buffer, "EDCB5433 ")); + + invoke_on_device(invokable::sprintf, buffer, "%0-20c", 'x'); + REQUIRE(!strcmp(buffer, "x ")); +} +#endif + +TEST_CASE("padding 20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%020d", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%020d", -1024); + REQUIRE(!strcmp(buffer, "-0000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%020i", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%020i", -1024); + REQUIRE(!strcmp(buffer, "-0000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%020u", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%020u", 4294966272U); + REQUIRE(!strcmp(buffer, "00000000004294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%020o", 511); + REQUIRE(!strcmp(buffer, "00000000000000000777")); + + invoke_on_device(invokable::sprintf, buffer, "%020o", 4294966785U); + REQUIRE(!strcmp(buffer, "00000000037777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%020x", 305441741); + REQUIRE(!strcmp(buffer, "0000000000001234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%020x", 3989525555U); + REQUIRE(!strcmp(buffer, "000000000000edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%020X", 305441741); + REQUIRE(!strcmp(buffer, "0000000000001234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%020X", 3989525555U); + REQUIRE(!strcmp(buffer, "000000000000EDCB5433")); +} + + +TEST_CASE("padding .20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%.20d", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%.20d", -1024); + REQUIRE(!strcmp(buffer, "-00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%.20i", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%.20i", -1024); + REQUIRE(!strcmp(buffer, "-00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%.20u", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%.20u", 4294966272U); + REQUIRE(!strcmp(buffer, "00000000004294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%.20o", 511); + REQUIRE(!strcmp(buffer, "00000000000000000777")); + + invoke_on_device(invokable::sprintf, buffer, "%.20o", 4294966785U); + REQUIRE(!strcmp(buffer, "00000000037777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%.20x", 305441741); + REQUIRE(!strcmp(buffer, "0000000000001234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%.20x", 3989525555U); + REQUIRE(!strcmp(buffer, "000000000000edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%.20X", 305441741); + REQUIRE(!strcmp(buffer, "0000000000001234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%.20X", 3989525555U); + REQUIRE(!strcmp(buffer, "000000000000EDCB5433")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("padding #020 - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%#020d", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%#020d", -1024); + REQUIRE(!strcmp(buffer, "-0000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%#020i", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%#020i", -1024); + REQUIRE(!strcmp(buffer, "-0000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%#020u", 1024); + REQUIRE(!strcmp(buffer, "00000000000000001024")); + + invoke_on_device(invokable::sprintf, buffer, "%#020u", 4294966272U); + REQUIRE(!strcmp(buffer, "00000000004294966272")); +} +#endif + +TEST_CASE("padding #020", "[]" ) { + char buffer[base_buffer_size]; + invoke_on_device(invokable::sprintf, buffer, "%#020o", 511); + REQUIRE(!strcmp(buffer, "00000000000000000777")); + + invoke_on_device(invokable::sprintf, buffer, "%#020o", 4294966785U); + REQUIRE(!strcmp(buffer, "00000000037777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%#020x", 305441741); + REQUIRE(!strcmp(buffer, "0x00000000001234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%#020x", 3989525555U); + REQUIRE(!strcmp(buffer, "0x0000000000edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%#020X", 305441741); + REQUIRE(!strcmp(buffer, "0X00000000001234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%#020X", 3989525555U); + REQUIRE(!strcmp(buffer, "0X0000000000EDCB5433")); +} + + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("padding #20 - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%#20d", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%#20d", -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%#20i", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%#20i", -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%#20u", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%#20u", 4294966272U); + REQUIRE(!strcmp(buffer, " 4294966272")); +} +#endif + +TEST_CASE("padding #20", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%#20o", 511); + REQUIRE(!strcmp(buffer, " 0777")); + + invoke_on_device(invokable::sprintf, buffer, "%#20o", 4294966785U); + REQUIRE(!strcmp(buffer, " 037777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%#20x", 305441741); + REQUIRE(!strcmp(buffer, " 0x1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%#20x", 3989525555U); + REQUIRE(!strcmp(buffer, " 0xedcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%#20X", 305441741); + REQUIRE(!strcmp(buffer, " 0X1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%#20X", 3989525555U); + REQUIRE(!strcmp(buffer, " 0XEDCB5433")); +} + + +TEST_CASE("padding 20.5", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%20.5d", 1024); + REQUIRE(!strcmp(buffer, " 01024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5d", -1024); + REQUIRE(!strcmp(buffer, " -01024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5i", 1024); + REQUIRE(!strcmp(buffer, " 01024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5i", -1024); + REQUIRE(!strcmp(buffer, " -01024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5u", 1024); + REQUIRE(!strcmp(buffer, " 01024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5u", 4294966272U); + REQUIRE(!strcmp(buffer, " 4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5o", 511); + REQUIRE(!strcmp(buffer, " 00777")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5o", 4294966785U); + REQUIRE(!strcmp(buffer, " 37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5x", 305441741); + REQUIRE(!strcmp(buffer, " 1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%20.10x", 3989525555U); + REQUIRE(!strcmp(buffer, " 00edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%20.5X", 305441741); + REQUIRE(!strcmp(buffer, " 1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%20.10X", 3989525555U); + REQUIRE(!strcmp(buffer, " 00EDCB5433")); +} + + +TEST_CASE("padding neg numbers", "[]" ) { + char buffer[base_buffer_size]; + + // space padding + invoke_on_device(invokable::sprintf, buffer, "% 1d", -5); + REQUIRE(!strcmp(buffer, "-5")); + + invoke_on_device(invokable::sprintf, buffer, "% 2d", -5); + REQUIRE(!strcmp(buffer, "-5")); + + invoke_on_device(invokable::sprintf, buffer, "% 3d", -5); + REQUIRE(!strcmp(buffer, " -5")); + + invoke_on_device(invokable::sprintf, buffer, "% 4d", -5); + REQUIRE(!strcmp(buffer, " -5")); + + // zero padding + invoke_on_device(invokable::sprintf, buffer, "%01d", -5); + REQUIRE(!strcmp(buffer, "-5")); + + invoke_on_device(invokable::sprintf, buffer, "%02d", -5); + REQUIRE(!strcmp(buffer, "-5")); + + invoke_on_device(invokable::sprintf, buffer, "%03d", -5); + REQUIRE(!strcmp(buffer, "-05")); + + invoke_on_device(invokable::sprintf, buffer, "%04d", -5); + REQUIRE(!strcmp(buffer, "-005")); +} + + +TEST_CASE("float padding neg numbers", "[]" ) { + char buffer[base_buffer_size]; + + // space padding + invoke_on_device(invokable::sprintf, buffer, "% 3.1f", -5.); + REQUIRE(!strcmp(buffer, "-5.0")); + + invoke_on_device(invokable::sprintf, buffer, "% 4.1f", -5.); + REQUIRE(!strcmp(buffer, "-5.0")); + + invoke_on_device(invokable::sprintf, buffer, "% 5.1f", -5.); + REQUIRE(!strcmp(buffer, " -5.0")); + +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + invoke_on_device(invokable::sprintf, buffer, "% 6.1g", -5.); + REQUIRE(!strcmp(buffer, " -5")); + + invoke_on_device(invokable::sprintf, buffer, "% 6.1e", -5.); + REQUIRE(!strcmp(buffer, "-5.0e+00")); + + invoke_on_device(invokable::sprintf, buffer, "% 10.1e", -5.); + REQUIRE(!strcmp(buffer, " -5.0e+00")); +#endif + + // zero padding + invoke_on_device(invokable::sprintf, buffer, "%03.1f", -5.); + REQUIRE(!strcmp(buffer, "-5.0")); + + invoke_on_device(invokable::sprintf, buffer, "%04.1f", -5.); + REQUIRE(!strcmp(buffer, "-5.0")); + + invoke_on_device(invokable::sprintf, buffer, "%05.1f", -5.); + REQUIRE(!strcmp(buffer, "-05.0")); + + // zero padding no decimal point + invoke_on_device(invokable::sprintf, buffer, "%01.0f", -5.); + REQUIRE(!strcmp(buffer, "-5")); + + invoke_on_device(invokable::sprintf, buffer, "%02.0f", -5.); + REQUIRE(!strcmp(buffer, "-5")); + + invoke_on_device(invokable::sprintf, buffer, "%03.0f", -5.); + REQUIRE(!strcmp(buffer, "-05")); + +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + invoke_on_device(invokable::sprintf, buffer, "%010.1e", -5.); + REQUIRE(!strcmp(buffer, "-005.0e+00")); + + invoke_on_device(invokable::sprintf, buffer, "%07.0E", -5.); + REQUIRE(!strcmp(buffer, "-05E+00")); + + invoke_on_device(invokable::sprintf, buffer, "%03.0g", -5.); + REQUIRE(!strcmp(buffer, "-05")); +#endif +} + +TEST_CASE("length", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%.0s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, "")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%.s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, "")); + + invoke_on_device(invokable::sprintf, buffer, "%20.s", mds("Hello testing")); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0d", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0d", -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.d", 0); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0i", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.i", -1024); + REQUIRE(!strcmp(buffer, " -1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.i", 0); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%20.u", 1024); + REQUIRE(!strcmp(buffer, " 1024")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0u", 4294966272U); + REQUIRE(!strcmp(buffer, " 4294966272")); + + invoke_on_device(invokable::sprintf, buffer, "%20.u", 0U); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%20.o", 511); + REQUIRE(!strcmp(buffer, " 777")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0o", 4294966785U); + REQUIRE(!strcmp(buffer, " 37777777001")); + + invoke_on_device(invokable::sprintf, buffer, "%20.o", 0U); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%20.x", 305441741); + REQUIRE(!strcmp(buffer, " 1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%50.x", 305441741); + REQUIRE(!strcmp(buffer, " 1234abcd")); + + invoke_on_device(invokable::sprintf, buffer, "%50.x%10.u", 305441741, 12345); + REQUIRE(!strcmp(buffer, " 1234abcd 12345")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0x", 3989525555U); + REQUIRE(!strcmp(buffer, " edcb5433")); + + invoke_on_device(invokable::sprintf, buffer, "%20.x", 0U); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%20.X", 305441741); + REQUIRE(!strcmp(buffer, " 1234ABCD")); + + invoke_on_device(invokable::sprintf, buffer, "%20.0X", 3989525555U); + REQUIRE(!strcmp(buffer, " EDCB5433")); + + invoke_on_device(invokable::sprintf, buffer, "%20.X", 0U); + REQUIRE(!strcmp(buffer, " ")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("length - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%02.0u", 0U); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%02.0d", 0); + REQUIRE(!strcmp(buffer, " ")); +} +#endif + + +TEST_CASE("float", "[]" ) { + char buffer[base_buffer_size]; + + // test special-case floats using math.h macros + invoke_on_device(invokable::sprintf, buffer, "%8f", (double) NAN); + REQUIRE(!strcmp(buffer, " nan")); + + invoke_on_device(invokable::sprintf, buffer, "%8f", (double) INFINITY); + REQUIRE(!strcmp(buffer, " inf")); + + invoke_on_device(invokable::sprintf, buffer, "%-8f", (double) -INFINITY); + REQUIRE(!strcmp(buffer, "-inf ")); + +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + invoke_on_device(invokable::sprintf, buffer, "%+8e", (double) INFINITY); + REQUIRE(!strcmp(buffer, " +inf")); +#endif + + invoke_on_device(invokable::sprintf, buffer, "%.4f", 3.1415354); + REQUIRE(!strcmp(buffer, "3.1415")); + + invoke_on_device(invokable::sprintf, buffer, "%.3f", 30343.1415354); + REQUIRE(!strcmp(buffer, "30343.142")); + + // switch from decimal to exponential representation + // + invoke_on_device(invokable::sprintf, buffer, "%.0f", (double) ((int64_t)1 * 1000 ) ); + if (PRINTF_MAX_FLOAT < 10e+2) { + REQUIRE(!strcmp(buffer, "10e+2")); + } + else { + REQUIRE(!strcmp(buffer, "1000")); + } + + invoke_on_device(invokable::sprintf, buffer, "%.0f", (double) ((int64_t)1 * 1000 * 1000 ) ); + if (PRINTF_MAX_FLOAT < 10e+5) { + REQUIRE(!strcmp(buffer, "10e+5")); + } + else { + REQUIRE(!strcmp(buffer, "1000000")); + } + + invoke_on_device(invokable::sprintf, buffer, "%.0f", (double) ((int64_t)1 * 1000 * 1000 * 1000 ) ); + if (PRINTF_MAX_FLOAT < 10e+8) { + REQUIRE(!strcmp(buffer, "10e+8")); + } + else { + REQUIRE(!strcmp(buffer, "1000000000")); + } + + invoke_on_device(invokable::sprintf, buffer, "%.0f", (double) ((int64_t)1 * 1000 * 1000 * 1000 * 1000) ); + if (PRINTF_MAX_FLOAT < 10e+11) { + REQUIRE(!strcmp(buffer, "10e+11")); + } + else { + REQUIRE(!strcmp(buffer, "1000000000000")); + } + + invoke_on_device(invokable::sprintf, buffer, "%.0f", (double) ((int64_t)1 * 1000 * 1000 * 1000 * 1000 * 1000) ); + if (PRINTF_MAX_FLOAT < 10e+14) { + REQUIRE(!strcmp(buffer, "10e+14")); + } + else { + REQUIRE(!strcmp(buffer, "1000000000000000")); + } + + invoke_on_device(invokable::sprintf, buffer, "%.0f", 34.1415354); + REQUIRE(!strcmp(buffer, "34")); + + invoke_on_device(invokable::sprintf, buffer, "%.0f", 1.3); + REQUIRE(!strcmp(buffer, "1")); + + invoke_on_device(invokable::sprintf, buffer, "%.0f", 1.55); + REQUIRE(!strcmp(buffer, "2")); + + invoke_on_device(invokable::sprintf, buffer, "%.1f", 1.64); + REQUIRE(!strcmp(buffer, "1.6")); + + invoke_on_device(invokable::sprintf, buffer, "%.2f", 42.8952); + REQUIRE(!strcmp(buffer, "42.90")); + + invoke_on_device(invokable::sprintf, buffer, "%.9f", 42.8952); + REQUIRE(!strcmp(buffer, "42.895200000")); + + invoke_on_device(invokable::sprintf, buffer, "%.10f", 42.895223); + REQUIRE(!strcmp(buffer, "42.8952230000")); + + // this testcase checks, that the precision is truncated to 9 digits. + // a perfect working float should return the whole number + invoke_on_device(invokable::sprintf, buffer, "%.12f", 42.89522312345678); + REQUIRE(!strcmp(buffer, "42.895223123000")); + + // this testcase checks, that the precision is truncated AND rounded to 9 digits. + // a perfect working float should return the whole number + invoke_on_device(invokable::sprintf, buffer, "%.12f", 42.89522387654321); + REQUIRE(!strcmp(buffer, "42.895223877000")); + + invoke_on_device(invokable::sprintf, buffer, "%6.2f", 42.8952); + REQUIRE(!strcmp(buffer, " 42.90")); + + invoke_on_device(invokable::sprintf, buffer, "%+6.2f", 42.8952); + REQUIRE(!strcmp(buffer, "+42.90")); + + invoke_on_device(invokable::sprintf, buffer, "%+5.1f", 42.9252); + REQUIRE(!strcmp(buffer, "+42.9")); + + invoke_on_device(invokable::sprintf, buffer, "%f", 42.5); + REQUIRE(!strcmp(buffer, "42.500000")); + + invoke_on_device(invokable::sprintf, buffer, "%.1f", 42.5); + REQUIRE(!strcmp(buffer, "42.5")); + + invoke_on_device(invokable::sprintf, buffer, "%f", 42167.0); + REQUIRE(!strcmp(buffer, "42167.000000")); + + invoke_on_device(invokable::sprintf, buffer, "%.9f", -12345.987654321); + REQUIRE(!strcmp(buffer, "-12345.987654321")); + + invoke_on_device(invokable::sprintf, buffer, "%.1f", 3.999); + REQUIRE(!strcmp(buffer, "4.0")); + + invoke_on_device(invokable::sprintf, buffer, "%.0f", 3.5); + REQUIRE(!strcmp(buffer, "4")); + + invoke_on_device(invokable::sprintf, buffer, "%.0f", 4.5); + REQUIRE(!strcmp(buffer, "4")); + + invoke_on_device(invokable::sprintf, buffer, "%.0f", 3.49); + REQUIRE(!strcmp(buffer, "3")); + + invoke_on_device(invokable::sprintf, buffer, "%.1f", 3.49); + REQUIRE(!strcmp(buffer, "3.5")); + + invoke_on_device(invokable::sprintf, buffer, "a%-5.1f", 0.5); + REQUIRE(!strcmp(buffer, "a0.5 ")); + + invoke_on_device(invokable::sprintf, buffer, "a%-5.1fend", 0.5); + REQUIRE(!strcmp(buffer, "a0.5 end")); + +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + invoke_on_device(invokable::sprintf, buffer, "%G", 12345.678); + REQUIRE(!strcmp(buffer, "12345.7")); + + invoke_on_device(invokable::sprintf, buffer, "%.7G", 12345.678); + REQUIRE(!strcmp(buffer, "12345.68")); + + invoke_on_device(invokable::sprintf, buffer, "%.5G", 123456789.); + REQUIRE(!strcmp(buffer, "1.2346E+08")); + + invoke_on_device(invokable::sprintf, buffer, "%.6G", 12345.); + REQUIRE(!strcmp(buffer, "12345.0")); + + invoke_on_device(invokable::sprintf, buffer, "%+12.4g", 123456789.); + REQUIRE(!strcmp(buffer, " +1.235e+08")); + + invoke_on_device(invokable::sprintf, buffer, "%.2G", 0.001234); + REQUIRE(!strcmp(buffer, "0.0012")); + + invoke_on_device(invokable::sprintf, buffer, "%+10.4G", 0.001234); + REQUIRE(!strcmp(buffer, " +0.001234")); + + invoke_on_device(invokable::sprintf, buffer, "%+012.4g", 0.00001234); + REQUIRE(!strcmp(buffer, "+001.234e-05")); + + invoke_on_device(invokable::sprintf, buffer, "%.3g", -1.2345e-308); + REQUIRE(!strcmp(buffer, "-1.23e-308")); + + invoke_on_device(invokable::sprintf, buffer, "%+.3E", 1.23e+308); + REQUIRE(!strcmp(buffer, "+1.230E+308")); +#endif + + // out of range for float: should switch to exp notation if supported, else empty + invoke_on_device(invokable::sprintf, buffer, "%.1f", 1E20); +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + REQUIRE(!strcmp(buffer, "1.0e+20")); +#else + REQUIRE(!strcmp(buffer, "")); +#endif + + // brute force float + bool fail = false; + std::stringstream str; + str.precision(5); + for (float i = -100000; i < 100000; i += 1) { + invoke_on_device(invokable::sprintf, buffer, "%.5f", (double)(i / 10000)); + str.str(""); + str << std::fixed << i / 10000; + fail = fail || !!strcmp(buffer, str.str().c_str()); + } + REQUIRE(!fail); + + +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + // brute force exp + str.setf(std::ios::scientific, std::ios::floatfield); + for (float i = -1e20; i < (float) 1e20; i += (float) 1e15) { + invoke_on_device(invokable::sprintf, buffer, "%.5f", (double) i); + str.str(""); + str << i; + fail = fail || !!strcmp(buffer, str.str().c_str()); + } + REQUIRE(!fail); +#endif +} + + +TEST_CASE("types", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%i", 0); + REQUIRE(!strcmp(buffer, "0")); + + invoke_on_device(invokable::sprintf, buffer, "%i", 1234); + REQUIRE(!strcmp(buffer, "1234")); + + invoke_on_device(invokable::sprintf, buffer, "%i", 32767); + REQUIRE(!strcmp(buffer, "32767")); + + invoke_on_device(invokable::sprintf, buffer, "%i", -32767); + REQUIRE(!strcmp(buffer, "-32767")); + + invoke_on_device(invokable::sprintf, buffer, "%li", 30L); + REQUIRE(!strcmp(buffer, "30")); + + invoke_on_device(invokable::sprintf, buffer, "%li", -2147483647L); + REQUIRE(!strcmp(buffer, "-2147483647")); + + invoke_on_device(invokable::sprintf, buffer, "%li", 2147483647L); + REQUIRE(!strcmp(buffer, "2147483647")); + + invoke_on_device(invokable::sprintf, buffer, "%lli", 30LL); + REQUIRE(!strcmp(buffer, "30")); + + invoke_on_device(invokable::sprintf, buffer, "%lli", -9223372036854775807LL); + REQUIRE(!strcmp(buffer, "-9223372036854775807")); + + invoke_on_device(invokable::sprintf, buffer, "%lli", 9223372036854775807LL); + REQUIRE(!strcmp(buffer, "9223372036854775807")); + + invoke_on_device(invokable::sprintf, buffer, "%lu", 100000L); + REQUIRE(!strcmp(buffer, "100000")); + + invoke_on_device(invokable::sprintf, buffer, "%lu", 0xFFFFFFFFL); + REQUIRE(!strcmp(buffer, "4294967295")); + + invoke_on_device(invokable::sprintf, buffer, "%llu", 281474976710656LLU); + REQUIRE(!strcmp(buffer, "281474976710656")); + + invoke_on_device(invokable::sprintf, buffer, "%llu", 18446744073709551615LLU); + REQUIRE(!strcmp(buffer, "18446744073709551615")); + + invoke_on_device(invokable::sprintf, buffer, "%zu", (size_t)2147483647UL); + REQUIRE(!strcmp(buffer, "2147483647")); + + invoke_on_device(invokable::sprintf, buffer, "%zd", (size_t)2147483647UL); + REQUIRE(!strcmp(buffer, "2147483647")); + + invoke_on_device(invokable::sprintf, buffer, "%zi", (ssize_t)-2147483647L); + REQUIRE(!strcmp(buffer, "-2147483647")); + + invoke_on_device(invokable::sprintf, buffer, "%o", 60000); + REQUIRE(!strcmp(buffer, "165140")); + + invoke_on_device(invokable::sprintf, buffer, "%lo", 12345678L); + REQUIRE(!strcmp(buffer, "57060516")); + + invoke_on_device(invokable::sprintf, buffer, "%lx", 0x12345678L); + REQUIRE(!strcmp(buffer, "12345678")); + + invoke_on_device(invokable::sprintf, buffer, "%llx", 0x1234567891234567LLU); + REQUIRE(!strcmp(buffer, "1234567891234567")); + + invoke_on_device(invokable::sprintf, buffer, "%lx", 0xabcdefabL); + REQUIRE(!strcmp(buffer, "abcdefab")); + + invoke_on_device(invokable::sprintf, buffer, "%lX", 0xabcdefabL); + REQUIRE(!strcmp(buffer, "ABCDEFAB")); + + invoke_on_device(invokable::sprintf, buffer, "%c", 'v'); + REQUIRE(!strcmp(buffer, "v")); + + invoke_on_device(invokable::sprintf, buffer, "%cv", 'w'); + REQUIRE(!strcmp(buffer, "wv")); + + invoke_on_device(invokable::sprintf, buffer, "%s", mds("A Test")); + REQUIRE(!strcmp(buffer, "A Test")); + + invoke_on_device(invokable::sprintf, buffer, "%hhu", (unsigned char) 0xFFU); + REQUIRE(!strcmp(buffer, "255")); + + invoke_on_device(invokable::sprintf, buffer, "%hu", (unsigned short) 0x1234u); + REQUIRE(!strcmp(buffer, "4660")); + + invoke_on_device(invokable::sprintf, buffer, "%s%hhi %hu", mds("Test"), (char) 100, (unsigned short) 0xFFFF); + REQUIRE(!strcmp(buffer, "Test100 65535")); + + invoke_on_device(invokable::sprintf, buffer, "%tx", &buffer[10] - &buffer[0]); + REQUIRE(!strcmp(buffer, "a")); + + invoke_on_device(invokable::sprintf, buffer, "%ji", (intmax_t)-2147483647L); + REQUIRE(!strcmp(buffer, "-2147483647")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("types - non-standard format", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%b", 60000); + REQUIRE(!strcmp(buffer, "1110101001100000")); + + invoke_on_device(invokable::sprintf, buffer, "%lb", 12345678L); + REQUIRE(!strcmp(buffer, "101111000110000101001110")); +} +#endif + +TEST_CASE("pointer", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%p", (void*)0x1234U); + if (sizeof(void*) == 4U) { + REQUIRE(!strcmp(buffer, "0x00001234")); + } + else { + REQUIRE(!strcmp(buffer, "0x0000000000001234")); + } + + invoke_on_device(invokable::sprintf, buffer, "%p", (void*)0x12345678U); + if (sizeof(void*) == 4U) { + REQUIRE(!strcmp(buffer, "0x12345678")); + } + else { + REQUIRE(!strcmp(buffer, "0x0000000012345678")); + } + + invoke_on_device(invokable::sprintf, buffer, "%p-%p", (void*)0x12345678U, (void*)0x7EDCBA98U); + if (sizeof(void*) == 4U) { + REQUIRE(!strcmp(buffer, "0x12345678-0x7edcba98")); + } + else { + REQUIRE(!strcmp(buffer, "0x0000000012345678-0x000000007edcba98")); + } + + if (sizeof(uintptr_t) == sizeof(uint64_t)) { + invoke_on_device(invokable::sprintf, buffer, "%p", (void*)(uintptr_t)0xFFFFFFFFU); + REQUIRE(!strcmp(buffer, "0x00000000ffffffff")); + } + else { + invoke_on_device(invokable::sprintf, buffer, "%p", (void*)(uintptr_t)0xFFFFFFFFU); + REQUIRE(!strcmp(buffer, "0xffffffff")); + } + + invoke_on_device(invokable::sprintf, buffer, "%p", NULL); + REQUIRE(!strcmp(buffer, "(nil)")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("unknown flag (non-standard format)", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%kmarco", 42, 37); + REQUIRE(!strcmp(buffer, "kmarco")); +} +#endif + +TEST_CASE("string length", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%.4s", mds("This is a test")); + REQUIRE(!strcmp(buffer, "This")); + + invoke_on_device(invokable::sprintf, buffer, "%.4s", mds("test")); + REQUIRE(!strcmp(buffer, "test")); + + invoke_on_device(invokable::sprintf, buffer, "%.7s", mds("123")); + REQUIRE(!strcmp(buffer, "123")); + + invoke_on_device(invokable::sprintf, buffer, "%.7s", mds("")); + REQUIRE(!strcmp(buffer, "")); + + invoke_on_device(invokable::sprintf, buffer, "%.4s%.2s", mds("123456"), mds("abcdef")); + REQUIRE(!strcmp(buffer, "1234ab")); + + invoke_on_device(invokable::sprintf, buffer, "%.*s", 3, mds("123456")); + REQUIRE(!strcmp(buffer, "123")); + +DISABLE_WARNING_PUSH +DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + invoke_on_device(invokable::sprintf, buffer, "%.*s", 3, NULL); +DISABLE_WARNING_POP + REQUIRE(!strcmp(buffer, "(null)")); +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +TEST_CASE("string length (non-standard format)", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%.4.2s", mds("123456")); + REQUIRE(!strcmp(buffer, ".2s")); +} +#endif + + +TEST_CASE("buffer length", "[]" ) { + char buffer[base_buffer_size]; + int ret; + + ret = invoke_on_device(invokable::snprintf, nullptr, 10, "%s", mds("Test")); + REQUIRE(ret == 4); + ret = invoke_on_device(invokable::snprintf, nullptr, (size_t) 0, "%s", mds("Test")); + REQUIRE(ret == 4); + + buffer[0] = (char)0xA5; + ret = invoke_on_device(invokable::snprintf, buffer, (size_t) 0, "%s", mds("Test")); + REQUIRE(buffer[0] == (char)0xA5); + REQUIRE(ret == 4); + + buffer[0] = (char)0xCC; + invoke_on_device(invokable::snprintf, buffer, 1, "%s", mds("Test")); + REQUIRE(buffer[0] == '\0'); + + invoke_on_device(invokable::snprintf, buffer, 2, "%s", mds("Hello")); + REQUIRE(!strcmp(buffer, "H")); + +DISABLE_WARNING_PUSH +DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + invoke_on_device(invokable::snprintf, buffer, 2, "%s", NULL); +DISABLE_WARNING_POP + REQUIRE(!strcmp(buffer, "(")); +} + + +TEST_CASE("ret value", "[]" ) { + char buffer[base_buffer_size]; + int ret; + + ret = invoke_on_device(invokable::snprintf, buffer, 6, "0%s", mds("1234")); + REQUIRE(!strcmp(buffer, "01234")); + REQUIRE(ret == 5); + + ret = invoke_on_device(invokable::snprintf, buffer, 6, "0%s", mds("12345")); + REQUIRE(!strcmp(buffer, "01234")); + REQUIRE(ret == 6); // "5" is truncated + + ret = invoke_on_device(invokable::snprintf, buffer, 6, "0%s", mds("1234567")); + REQUIRE(!strcmp(buffer, "01234")); + REQUIRE(ret == 8); // "567" are truncated + +DISABLE_WARNING_PUSH +DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + ret = invoke_on_device(invokable::snprintf, buffer, 6, "0%s", NULL); +DISABLE_WARNING_POP + REQUIRE(!strcmp(buffer, "0(nul")); + REQUIRE(ret == 7); // "l)" is truncated + + ret = invoke_on_device(invokable::snprintf, buffer, 10, "hello, world"); + REQUIRE(ret == 12); + + ret = invoke_on_device(invokable::snprintf, buffer, 3, "%d", 10000); + REQUIRE(ret == 5); + REQUIRE(strlen(buffer) == 2U); + REQUIRE(buffer[0] == '1'); + REQUIRE(buffer[1] == '0'); + REQUIRE(buffer[2] == '\0'); +} + + +TEST_CASE("misc", "[]" ) { + char buffer[base_buffer_size]; + + invoke_on_device(invokable::sprintf, buffer, "%u%u%ctest%d %s", 5, 3000, 'a', -20, mds("bit")); + REQUIRE(!strcmp(buffer, "53000atest-20 bit")); + + invoke_on_device(invokable::sprintf, buffer, "%.*f", 2, 0.33333333); + REQUIRE(!strcmp(buffer, "0.33")); + + invoke_on_device(invokable::sprintf, buffer, "%.*d", -1, 1); + REQUIRE(!strcmp(buffer, "1")); + + invoke_on_device(invokable::sprintf, buffer, "%.3s", mds("foobar")); + REQUIRE(!strcmp(buffer, "foo")); + + invoke_on_device(invokable::sprintf, buffer, "% .0d", 0); + REQUIRE(!strcmp(buffer, " ")); + + invoke_on_device(invokable::sprintf, buffer, "%10.5d", 4); + REQUIRE(!strcmp(buffer, " 00004")); + + invoke_on_device(invokable::sprintf, buffer, "%*sx", -3, mds("hi")); + REQUIRE(!strcmp(buffer, "hi x")); + +#ifndef PRINTF_DISABLE_SUPPORT_EXPONENTIAL + invoke_on_device(invokable::sprintf, buffer, "%.*g", 2, 0.33333333); + REQUIRE(!strcmp(buffer, "0.33")); + + invoke_on_device(invokable::sprintf, buffer, "%.*e", 2, 0.33333333); + REQUIRE(!strcmp(buffer, "3.33e-01")); +#endif +} + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +DISABLE_WARNING_POP +#endif + diff --git a/test/cuda_test_suite_host.cu b/test/cuda_test_suite_host.cu new file mode 100644 index 00000000..6390d430 --- /dev/null +++ b/test/cuda_test_suite_host.cu @@ -0,0 +1 @@ +#include "test_suite.cpp" diff --git a/test/test_suite.cpp b/test/test_suite.cpp index 86fe1ba3..49ee6492 100644 --- a/test/test_suite.cpp +++ b/test/test_suite.cpp @@ -50,8 +50,19 @@ namespace test { #define DISABLE_WARNING(warningNumber) __pragma(warning( disable : warningNumber )) // TODO: find the right warning number for this - #define DISABLE_WARNING_PRINTF_FORMAT - #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + +#elif defined(__NVCC__) + #define DO_PRAGMA(X) _Pragma(#X) + #define DISABLE_WARNING_PUSH DO_PRAGMA(push) + #define DISABLE_WARNING_POP DO_PRAGMA(pop) + #define DISABLE_WARNING(warning_code) DO_PRAGMA(diag_suppress warning_code) + + #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(bad_printf_format_string) + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW #elif defined(__GNUC__) || defined(__clang__) #define DO_PRAGMA(X) _Pragma(#X) @@ -59,15 +70,15 @@ namespace test { #define DISABLE_WARNING_POP DO_PRAGMA(GCC diagnostic pop) #define DISABLE_WARNING(warningName) DO_PRAGMA(GCC diagnostic ignored #warningName) - #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(-Wformat) - #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS DISABLE_WARNING(-Wformat-extra-args) - #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW DISABLE_WARNING(-Wformat-overflow) - + #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(-Wformat) + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS DISABLE_WARNING(-Wformat-extra-args) + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW DISABLE_WARNING(-Wformat-overflow) #else #define DISABLE_WARNING_PUSH #define DISABLE_WARNING_POP #define DISABLE_WARNING_PRINTF_FORMAT #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW #endif #ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS