From f4aa2b99774f4cc9b4967cac3eacdba57ad5c63e Mon Sep 17 00:00:00 2001 From: TomSang Date: Tue, 14 Sep 2021 04:06:32 -0400 Subject: [PATCH] SWDEV-266829 - Enable more tests on AMD and NV devices (#2336) * SWDEV-266829 - Enable more tests on AMD and NV devices 1. Enable tests on AMD and NV devices tests/src/runtimeApi/event/hipEventMultiThreaded.cpp Loops and threads per core are changed smaller so that test can finish in a shorter time. tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp Fix logic error on how to get priority_normal 2. Simply enable test on AMD device tests/src/runtimeApi/memory/hipManagedKeyword.cpp tests/src/runtimeApi/module/hipManagedKeyword.cpp tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp tests/src/runtimeApi/memory/p2p_copy_coherency.cpp 3. Simply enable test on NV device tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp 4. Fix typo tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp 5. Remove useless tests tests/src/hipC.c tests/src/hipHcc.cpp Change-Id: Ia4406353e64d69bd34c58ebb56185701f7ce1caa * Remove tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp for cuda test Co-authored-by: anusha GodavarthySurya Co-authored-by: Jenkins [ROCm/hip commit: 3fd16c0b5be5abc3483e2edc0a77a25269cc3bc5] --- projects/hip/tests/src/hipC.c | 55 ------------------- projects/hip/tests/src/hipHcc.cpp | 52 ------------------ .../event/hipEventMultiThreaded.cpp | 13 +++-- .../runtimeApi/memory/hipManagedKeyword.cpp | 2 +- .../runtimeApi/memory/p2p_copy_coherency.cpp | 2 +- .../runtimeApi/module/hipManagedKeyword.cpp | 4 +- .../stream/hipStreamACb_MultiThread.cpp | 2 +- .../stream/hipStreamAddCallbackCatch.cpp | 2 +- .../stream/hipStreamCreateWithPriority.cpp | 15 +++-- 9 files changed, 24 insertions(+), 123 deletions(-) delete mode 100644 projects/hip/tests/src/hipC.c delete mode 100644 projects/hip/tests/src/hipHcc.cpp diff --git a/projects/hip/tests/src/hipC.c b/projects/hip/tests/src/hipC.c deleted file mode 100644 index a4f24eac8e..0000000000 --- a/projects/hip/tests/src/hipC.c +++ /dev/null @@ -1,55 +0,0 @@ -/* -Copyright (c) 2015 - 2021 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. -*/ - -/* HIT_START - * BUILD: %t %s EXCLUDE_HIP_PLATFORM all - * TEST: %t - * HIT_END - */ - -#include "hip/hip_runtime.h" -#include "test_common.h" -#include - -#define ITER 1<<20 -#define SIZE 1024*1024*sizeof(int) - -__global__ void Iter(int *Ad){ - int tx = threadIdx.x + blockIdx.x * blockDim.x; - if(tx == 0){ - for(int i=0;i -#include -#include "hip/hip_runtime.h" -#include "hip/hip_ext.h" -#include "test_common.h" - -#define CHECK(error) \ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, __FILE__, \ - __LINE__); \ - exit(EXIT_FAILURE); \ - } - - -int main(int argc, char* argv[]) { - int deviceId; - CHECK(hipGetDevice(&deviceId)); - hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, deviceId)); - printf("info: running on device #%d %s\n", deviceId, props.name); - - passed(); -}; diff --git a/projects/hip/tests/src/runtimeApi/event/hipEventMultiThreaded.cpp b/projects/hip/tests/src/runtimeApi/event/hipEventMultiThreaded.cpp index 5c727d42d4..65f2ff1851 100644 --- a/projects/hip/tests/src/runtimeApi/event/hipEventMultiThreaded.cpp +++ b/projects/hip/tests/src/runtimeApi/event/hipEventMultiThreaded.cpp @@ -21,8 +21,8 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia - * TEST: %t EXCLUDE_HIP_PLATFORM amd + * BUILD: %t %s ../../test_common.cpp + * TEST: %t * HIT_END */ @@ -30,9 +30,12 @@ THE SOFTWARE. #include #include -#define THREADS 8 +#define THREADS 2 // threads per core #define MAX_NUM_THREADS 512 -#define ITER 100 +#define ITER 5 // total loop number + +// 5 loops and 2 threads per core are enough for function verification. +// You may adjust them for your test purpose. extern "C" __global__ void WaitKernel(int *Ad, int clockrate) { uint64_t wait_t = 500, @@ -73,7 +76,9 @@ int main(int argc, char* argv[]) { hipEventCreate(&start); std::thread t[NUM_THREADS]; + printf("NUM_THREADS=%d\n", NUM_THREADS); for (int i = 0; i < ITER; i++) { + printf("loop %d/%d\n", i, ITER); for (int j = 0; j < NUM_THREADS; j++) { t[j] = std::thread(t1, start, stream1, clkRate, A[j], Ad[j]); } diff --git a/projects/hip/tests/src/runtimeApi/memory/hipManagedKeyword.cpp b/projects/hip/tests/src/runtimeApi/memory/hipManagedKeyword.cpp index 5cdfbfa9f3..33b114c97b 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipManagedKeyword.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipManagedKeyword.cpp @@ -4,7 +4,7 @@ //Enable test when compiler support is available in mainline /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM amd + * BUILD: %t %s ../../test_common.cpp * HIT_END */ #define N 1048576 diff --git a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp index 056a0d7b03..2050bd38a6 100644 --- a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 - * TEST: %t EXCLUDE_HIP_PLATFORM all + * TEST: %t * HIT_END */ diff --git a/projects/hip/tests/src/runtimeApi/module/hipManagedKeyword.cpp b/projects/hip/tests/src/runtimeApi/module/hipManagedKeyword.cpp index 9bc6bc3e54..6e378e14e7 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipManagedKeyword.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipManagedKeyword.cpp @@ -21,8 +21,8 @@ THE SOFTWARE. */ /* HIT_START - * BUILD_CMD: managed_kernel.code %hc --genco %S/managed_kernel.cpp -o managed_kernel.code EXCLUDE_HIP_PLATFORM amd - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM amd + * BUILD_CMD: managed_kernel.code %hc --genco %S/managed_kernel.cpp -o managed_kernel.code + * BUILD: %t %s ../../test_common.cpp * TEST: %t * HIT_END */ diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp index edab4f2977..d26e41019c 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. // This test case is disabled currently. /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 * TEST: %t * HIT_END */ diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp index 4a0c419a88..c694bf6d2a 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp @@ -11,7 +11,7 @@ #include "test_common.h" /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * TEST: %t * HIT_END */ diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp index 782258fefa..af9f4497c6 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../test_common.cpp * TEST: %t * HIT_END */ @@ -55,7 +55,7 @@ __global__ void memcpy_kernel(T* dst, T* src, size_t n) } template -void runTest() +int runTest() { size_t size = NUMITERS*MEMCPYSIZE; @@ -71,13 +71,16 @@ void runTest() printf("HIP stream priority range - low: %d to high: %d\n", priority_low, priority_high); // Check if priorities are indeed supported - if ((priority_low - priority_high) == 0) { passed(); } - + if (priority_low == 0 && priority_high == 0) { + printf("The device doesn't support stream priorities\n"); + passed(); + } + // Enable/disable priorities based on number of available priority levels enable_priority_low = true; enable_priority_high = true; if ((priority_low - priority_high) > 1) enable_priority_normal = true; - if (enable_priority_normal) priority_normal = ((priority_low - priority_high) / 2); + if (enable_priority_normal) priority_normal = ((priority_low + priority_high) / 2); // create streams with highest and lowest available priorities #define OP(x) \ @@ -216,5 +219,5 @@ void runTest() int main(int argc, char **argv) { HipTest::parseStandardArguments(argc, argv, false); - runTest(); + return runTest(); }