diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index eb7b91d..9415e94 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1,15 +1,61 @@ -cmake_minimum_required(VERSION 3.9) +cmake_minimum_required(VERSION 3.20) enable_language(CXX) +option(BUILD_CUDA_TESTS "Include tests for use of the printf library with CUDA host-side and device-side code" OFF) + set(test_targets autotest) if (NOT ALIAS_STANDARD_FUNCTION_NAMES) list(APPEND test_targets test_suite) + set(targets_needing_config_h test_suite) endif() + option(TEST_WITH_NON_STANDARD_FORMAT_STRINGS "Include tests using non-standard-compliant format strings?" ON) # ... don't worry, we'll suppress the compiler warnings for those. +if(BUILD_CUDA_TESTS) + enable_language(CUDA) + list(APPEND CMAKE_CUDA_FLAGS "--extended-lambda --expt-relaxed-constexpr -Xcudafe --display_error_number") + cmake_policy(SET CMP0104 OLD) + include(FindCUDA/select_compute_arch) + list(APPEND cuda_test_targets 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}") + + foreach(tgt ${cuda_test_targets}) + string(REPLACE "cuda_" "cuda/" source_file_prefix ${tgt}) + add_executable(${tgt} "${source_file_prefix}.cu") + target_include_directories(${tgt} PRIVATE "$") + set_target_properties( + ${tgt} + PROPERTIES + CUDA_STANDARD 11 + CUDA_STANDARD_REQUIRED YES + CUDA_EXTENSIONS NO + ) + if (TEST_WITH_NON_STANDARD_FORMAT_STRINGS) + target_compile_definitions(${tgt} PRIVATE TEST_WITH_NON_STANDARD_FORMAT_STRINGS) + endif() + add_test(NAME "${PROJECT_NAME}.${tgt}" COMMAND "${tgt}") + endforeach() +endif() + +if(BUILD_CUDA_TESTS) + if (NOT ALIAS_STANDARD_FUNCTION_NAMES) + target_link_libraries(cuda_test_suite_device printf_cuda) + set_target_properties(cuda_test_suite_device PROPERTIES CUDA_SEPARABLE_COMPILATION YES) + list(APPEND targets_needing_config_h cuda_test_suite_host) +# target_compile_options(cuda_test_suite_device PRIVATE -G -g) + endif() +endif() + foreach(tgt ${test_targets}) add_executable(${tgt} "${tgt}.cpp") set_target_properties( @@ -75,20 +121,21 @@ foreach(tgt ${test_targets}) target_compile_options(${tgt} PRIVATE -ffat-lto-objects) endif() endif() - endforeach() if (NOT ALIAS_STANDARD_FUNCTION_NAMES) - # These two lines are necessary, since the test suite does not actually use the - # compiled library - it includes the library's source .c file; and that means we - # need to include the generated config.h file. - target_compile_definitions(test_suite PRIVATE PRINTF_INCLUDE_CONFIG_H) - target_include_directories( - test_suite - PRIVATE - "${GENERATED_INCLUDE_DIR}" - "$" - ) + # These following is necessary, since thee targets applying printf_config.h + # do not actually use the compiled library - they includes the library's source .c file; + # so we need to make sure it's accessible to them + foreach(tgt ${targets_needing_config_h}) + target_compile_definitions(test_suite PRIVATE PRINTF_INCLUDE_CONFIG_H) + target_include_directories( + ${tgt} + PRIVATE + "${GENERATED_INCLUDE_DIR}" + "$" + ) + endforeach() 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 0000000..53f14c8 --- /dev/null +++ b/test/cuda/test_suite_device.cu @@ -0,0 +1,553 @@ +#include "printf_config.h" +#include + +#include +#include +#include +#include +#include +#include + +// 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 + #define DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER + +#elif defined(__NVCC__) + #define DO_PRAGMA(X) _Pragma(#X) + #define DISABLE_WARNING_PUSH DO_PRAGMA(push) + #define DISABLE_WARNING_POP DO_PRAGMA(pop) + #ifdef __NVCC_DIAG_PRAGMA_SUPPORT__ + #define DISABLE_WARNING(warning_code) DO_PRAGMA(nv_diag_suppress warning_code) + #else + #define DISABLE_WARNING(warning_code) DO_PRAGMA(diag_suppress warning_code) + #endif + + #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(bad_printf_format_string) + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + #define DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER + + +#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) +#if defined(__clang__) + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + #define DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER DISABLE_WARNING(-Wformat-invalid-specifier) +#else + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW DISABLE_WARNING(-Wformat-overflow) + #define DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER +#endif +#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_INVALID_SPECIFIER +#endif + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +DISABLE_WARNING_PUSH +DISABLE_WARNING_PRINTF_FORMAT +DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS +DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER +#endif + +#if defined(_MSC_VER) +DISABLE_WARNING(4996) // Discouragement of use of std::sprintf() +DISABLE_WARNING(4310) // Casting to smaller type +DISABLE_WARNING(4127) // Constant conditional expression +#endif + +char* make_device_string(char const* s) +{ + if (s == nullptr) { + return nullptr; + } + + // Maybe it's _already_ a device string? + + cudaPointerAttributes attrs; + auto error = cudaPointerGetAttributes (&attrs, s); + if (error != cudaSuccess) { + throw std::runtime_error("CUDA error: " + std::string(cudaGetErrorString(error))); + } + switch(attrs.type) { + case cudaMemoryTypeUnregistered: // host mem, not registered with CUDA + case cudaMemoryTypeHost : break; + case cudaMemoryTypeDevice : + case cudaMemoryTypeManaged : + throw std::invalid_argument("Got a pointer which is already GPU device-side accessible"); + default: + throw std::invalid_argument("Get a pointer to an unsupported/unregistered memory type"); + } + + size_t size = strlen(s) + 1; + void* dsptr; + cudaMalloc(&dsptr, size); + cudaMemcpy(dsptr, s, size, cudaMemcpyDefault); + cudaDeviceSynchronize(); + return (char *) dsptr; +} + +inline char* mkstr(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 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 = 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 = 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 = sprintf_(buffer, format, args...); break; + case invokable::snprintf_: *result = 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; + 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 or buffer_size == 0) { + cudaMalloc(&buffer_d, buffer_size); + cudaMemcpy(buffer_d, buffer, buffer_size, cudaMemcpyDefault); + } else { + buffer_d = nullptr; + } + char* format_d = make_device_string(format); + // std::cout << "Copying done, now launching kernel." << std::endl; + 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...); +} + +template +int printing_check( + const char *expected, + const char *, + invokable invokable_printer, + char *buffer, + size_t buffer_size, + const char *format, + Ts &&... params) +{ + if (buffer == nullptr and expected != nullptr) { + std::cerr << "Internal error: A null buffer is expected to become non-null" << std::endl; + exit(EXIT_FAILURE); + } + auto ret = invoke_on_device(invokable_printer, buffer, buffer_size, format, std::forward(params)...); + // std::cout << "invoked_on_device with format \"" << format << "\" done." << std::endl; + if (buffer == nullptr) { + return ret; + } + if (buffer_size != base_buffer_size) { + buffer[base_buffer_size - 1] = '\0'; + } + // std::cout << "----\n"; + // std::cout << "Resulting buffer contents: " << '"' << buffer << '"' << '\n'; + if (strncmp(buffer, expected, buffer_size) != 0) { + buffer[strlen(expected)] = '\0'; + std::cerr << "Failed with printer " << name(invokable_printer) << + " with format \"" << format << "\":\n" + << "Actual: \"" << buffer << "\"\n" + << "Expected: \"" << expected << "\"\n" << std::flush; + exit(EXIT_FAILURE); + } + return ret; +} + +template +void printing_and_ret_check( + int expected_return_value, + const char *expected, + const char *, + invokable invokable_printer, + char *buffer, + size_t buffer_size, + const char *format, + Ts &&... params) +{ + auto ret = printing_check(expected, nullptr, invokable_printer, buffer, buffer_size, format, std::forward(params)...); + if (ret != expected_return_value) { + std::cerr << "Unexpected return value with printer " << name(invokable_printer) << + " and format \"" << format << "\":\n Actual: " << ret << "\n Expected: " << + expected_return_value << std::endl; + exit(EXIT_FAILURE); + } +} + +namespace kernels { + +__global__ void fctprintf_kernel(char* buffer) +{ + sv_and_pos svnp { {buffer, base_buffer_size}, 0 }; + fctprintf(append_to_buffer, &svnp, "This is a test of %X", 0x12EFU); +} + +} // namespace kernels + +void testcase_fctprintf() { + char buffer[base_buffer_size]; + char* buffer_d; + cudaMalloc(&buffer_d, base_buffer_size); + cudaMemset(buffer_d, 0xCC, base_buffer_size); + kernels::fctprintf_kernel<<<1, 1>>>(buffer_d); + cudaMemcpy(buffer, buffer_d, base_buffer_size, cudaMemcpyDefault); + cudaDeviceSynchronize(); + if (strncmp(buffer, "This is a test of 12EF", 22U) != 0) { + std::cerr << "fctprintf failed to produce the correct string." << std::endl; + exit(EXIT_FAILURE); + } + // Remember: printf does not append a `\0` to the output after going through its format string. + if (buffer[22] != (char)0xCC) { + std::cerr << "fctprintf changed buffer characters past where it was allowed to\n" << std::endl; + exit(EXIT_FAILURE); + } + 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); + 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 + +void testcase_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(); + if (strncmp(buffer, "This is a test of 12EF", 22U) != 0) { + std::cerr << "vfctprintf failed to produce the correct string." << std::endl; + exit(EXIT_FAILURE); + } + if (buffer[22] != (char)0xCC) { + std::cerr << "vfctprintf changed buffer characters past where it was allowed to\n" << std::endl; + exit(EXIT_FAILURE); + } + cudaFree(buffer_d); +} + +#define STRINGIFY(_x) #_x + +// You can add the following into PRINTING_CHECK_WITH_BUF_SIZE to keep +// track of which checks pass before a failure: +// +// printf("printing check at file %s line %d\n", __FILE__, (int) __LINE__); \ + +#define PRINTING_CHECK_WITH_BUF_SIZE(expected,dummy,printer,buffer,buffer_size,format,...) \ +do { \ + printing_check(expected, STRINGIFY(dummy), invokable::printer, buffer, buffer_size, format, ##__VA_ARGS__); \ +} while(false); + +#define PRINTING_CHECK(expected,dummy,printer,buffer,format,...) \ +PRINTING_CHECK_WITH_BUF_SIZE(expected, dummy, printer, buffer, base_buffer_size, format, ##__VA_ARGS__) + +void testcase_snprintf() { + char buffer[base_buffer_size]; + PRINTING_CHECK("-1000", ==, snprintf_, buffer, "%d", -1000); + PRINTING_CHECK_WITH_BUF_SIZE("-1", ==, snprintf_, buffer, 3, "%d", -1000); +} + +void testcase_vsprintf() { + char buffer[base_buffer_size]; + PRINTING_CHECK("-1", ==, vsprintf_, buffer, "%d", -1 ); + PRINTING_CHECK("3 -1000 test", ==, vsprintf_, buffer, "%d %d %s", 3, -1000, mkstr("test") ); +} + +void testcase_vsnprintf() { + char buffer[base_buffer_size]; + PRINTING_CHECK("-1", ==, vsnprintf_, buffer, "%d", -1); + PRINTING_CHECK("3 -1000 test", ==, vsnprintf_, buffer, "%d %d %s", 3, -1000, mkstr("test")); +} + +void testcase_simple_sprintf() { + char buffer[base_buffer_size]; + memset(buffer, 0xCC, base_buffer_size); + PRINTING_CHECK("42", ==, sprintf_, buffer, "%d", 42); +} + + +[[maybe_unused]] void testcase_brute_force_float() { + char buffer[base_buffer_size]; + + // 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()); + } + if (fail) { + std::cerr << "sprintf(\"" << "%.5f\" (double)(i / 10000)) failed." << std::endl; + exit(EXIT_FAILURE); + } + + +#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()); + } + if (fail) { + std::cerr << "sprintf(\"" << "%.5f\" (double) i) failed." << std::endl; + exit(EXIT_FAILURE); + } +#endif +} + + +#ifndef STRINGIFY +#define STRINGIFY(_x) #_x +#endif + +#define PRINTF_TEST_CASE(unstringified_name) void testcase_ ## unstringified_name() +#define CHECK(...) \ +do { \ + if (!(__VA_ARGS__)) { \ + std::cerr << "Check failed at " << __FILE__ << ':' << __LINE__ << std::endl; \ + exit(EXIT_FAILURE); \ + } \ +} while(0) + +#define CAPTURE_AND_PRINT(printer_, _buffer, format, ...) \ +do { \ + auto ret = invoke_on_device(invokable::printer_, _buffer, format, __VA_ARGS__); \ + if (not ret) { \ + std::cerr << "Failed at " << __FILE__ << ':' << __LINE__ << std::endl; \ + exit(EXIT_FAILURE); \ + } \ +} while(false); + +#define CAPTURE_AND_PRINT_WITH_RETVAL(retval, printer_, ...) \ +do { \ + retval = invoke_on_device(invokable::printer_, __VA_ARGS__); \ +} while(false); + + +#include "../test_suite_main_testcases.hpp" + +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS +DISABLE_WARNING_POP +#endif + +int main() +{ + // testcases defined within this .cu file + + testcase_fctprintf(); + testcase_vfctprintf(); + testcase_snprintf(); + testcase_vsprintf(); + testcase_vsnprintf(); + testcase_simple_sprintf(); +// Too long and costly for running on the GPU... +// testcase_brute_force_float(); + + // testcases defined in the main testcases header + + testcase_space_flag(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_space_flag__non_standard_format(); +#endif + testcase_plus_flag(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_plus_flag__non_standard_format(); +#endif + testcase_zero_flag(); + testcase_minus_flag(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_minus_flag_and_non_standard_zero_modifier_for_integers(); +#endif + testcase_sharp_flag(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_sharp_flag__non_standard_format(); +#endif +#if PRINTF_SUPPORT_LONG_LONG + testcase_sharp_flag_with_long_long(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_sharp_flag_with_long_long__non_standard_format(); +#endif +#endif // PRINTF_SUPPORT_LONG_LONG + testcase_specifier(); + testcase_width(); + testcase_width_20(); + testcase_width_asterisk_20(); + testcase_width_minus_20(); + testcase_width_0_minus_20(); + testcase_padding_20(); + testcase_padding_dot_20(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_padding_sharp_020__non_standard_format(); +#endif + testcase_padding_sharp_020(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_padding_sharp_20__non_standard_format(); +#endif + testcase_padding_sharp_20(); + testcase_padding_20_point_5(); + testcase_padding_negative_numbers(); +#if PRINTF_SUPPORT_DECIMAL_SPECIFIERS || PRINTF_SUPPORT_EXPONENTIAL_SPECIFIERS + testcase_float_padding_negative_numbers(); +#endif + testcase_length(); + testcase_length__non_standard_format(); +#if PRINTF_SUPPORT_DECIMAL_SPECIFIERS || PRINTF_SUPPORT_EXPONENTIAL_SPECIFIERS + testcase_infinity_and_not_a_number_values(); + testcase_floating_point_specifiers_with_31_to_32_bit_integer_values(); + testcase_tiny_floating_point_values(); + testcase_fallback_from_decimal_to_exponential(); + testcase_floating_point_specifiers_precision_and_flags(); +#endif + testcase_integer_types(); + testcase_types__non_standard_format(); + testcase_pointer(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_unknown_flag__non_standard_format(); +#endif + testcase_string_length(); +#ifdef TEST_WITH_NON_STANDARD_FORMAT_STRINGS + testcase_string_length__non_standard_format(); +#endif + testcase_buffer_length(); + testcase_misc(); + testcase_extremal_signed_integer_values(); + testcase_extremal_unsigned_integer_values(); +} + diff --git a/test/cuda/test_suite_host.cu b/test/cuda/test_suite_host.cu new file mode 100644 index 0000000..9fef612 --- /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 92d7cb0..33a8265 100644 --- a/test/test_suite.cpp +++ b/test/test_suite.cpp @@ -117,6 +117,22 @@ do { \ #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW #define DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER +#elif defined(__NVCC__) + #define DO_PRAGMA(X) _Pragma(#X) + #define DISABLE_WARNING_PUSH DO_PRAGMA(push) + #define DISABLE_WARNING_POP DO_PRAGMA(pop) + #ifdef __NVCC_DIAG_PRAGMA_SUPPORT__ + #define DISABLE_WARNING(warning_code) DO_PRAGMA(nv_diag_suppress warning_code) + #else + #define DISABLE_WARNING(warning_code) DO_PRAGMA(diag_suppress warning_code) + #endif + + #define DISABLE_WARNING_PRINTF_FORMAT DISABLE_WARNING(bad_printf_format_string) + #define DISABLE_WARNING_PRINTF_FORMAT_EXTRA_ARGS + #define DISABLE_WARNING_PRINTF_FORMAT_OVERFLOW + #define DISABLE_WARNING_PRINTF_FORMAT_INVALID_SPECIFIER + + #elif defined(__GNUC__) || defined(__clang__) #define DO_PRAGMA(X) _Pragma(#X) #define DISABLE_WARNING_PUSH DO_PRAGMA(GCC diagnostic push)