EXSWHTEC-307 - Implement tests for printf and __hip_hc_<add/sub/mul>8pk device functions (#274)
Change-Id: I4669d3f0f8cc062e75cd2b0a8aeee51868dd5b09
This commit is contained in:
Nives Vukovic
committed by
Rakesh Roy
Rakesh Roy
orang tua
30f427277b
melakukan
08a41abaad
@@ -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"
|
||||
]
|
||||
|
||||
@@ -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"
|
||||
]
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
|
||||
/**
|
||||
* @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<char4> 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<char4> 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<char4> 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_<add/sub/mul>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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
|
||||
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)
|
||||
@@ -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); }
|
||||
)"};
|
||||
@@ -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)
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_process.hh>
|
||||
|
||||
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);
|
||||
|
||||
@@ -21,7 +21,7 @@ THE SOFTWARE.
|
||||
#include <hip_test_defgroups.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup printf
|
||||
* @addtogroup printf printf
|
||||
* @{
|
||||
* @ingroup PrintfTest
|
||||
* `int printf()` -
|
||||
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
#include <hip_test_context.hh>
|
||||
|
||||
__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() {
|
||||
|
||||
@@ -26,7 +26,7 @@ __global__ void run_printf(int *count) {
|
||||
*count = printf("Hello World");
|
||||
}
|
||||
/**
|
||||
* @addtogroup printf
|
||||
* @addtogroup printf printf
|
||||
* @{
|
||||
* @ingroup PrintfTest
|
||||
* `int printf()` -
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_process.hh>
|
||||
|
||||
/**
|
||||
* @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);
|
||||
}
|
||||
@@ -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 <hip_test_context.hh>
|
||||
|
||||
__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<unsigned char>(42), static_cast<unsigned char>(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<void>(hipDeviceSynchronize());
|
||||
}
|
||||
@@ -45,7 +45,7 @@ __global__ void kernel_printf_thread(int *count) {
|
||||
}
|
||||
|
||||
/**
|
||||
* @addtogroup printf
|
||||
* @addtogroup printf printf
|
||||
* @{
|
||||
* @ingroup PrintfTest
|
||||
* `int printf()` -
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_process.hh>
|
||||
|
||||
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);
|
||||
}
|
||||
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
#include <hip_test_context.hh>
|
||||
|
||||
__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 %.*, %*.*
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
|
||||
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); }
|
||||
@@ -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); }
|
||||
)"};
|
||||
Reference in New Issue
Block a user