EXSWHTEC-307 - Implement tests for printf and __hip_hc_<add/sub/mul>8pk device functions (#274)
Change-Id: I4669d3f0f8cc062e75cd2b0a8aeee51868dd5b09
[ROCm/hip-tests commit: 08a41abaad]
此提交包含在：
Nives Vukovic
Rakesh Roy
@@ -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); }
|
||||
)"};
|
||||
新增問題並參考
封鎖使用者