AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

SWDEV-266829 - Enable more tests on AMD and NV devices (#2336)

Quellcode durchsuchen 
* 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 <Anusha.GodavarthySurya@amd.com>
Co-authored-by: Jenkins <jenkins-compute@amd.com>

[ROCm/hip commit: 3fd16c0b5b]
Dieser Commit ist enthalten in:
TomSang
2021-09-14 04:06:32 -04:00
committet von GitHub
Ursprung 12c15c7f71
Commit f4aa2b9977
9 geänderte Dateien mit 24 neuen und 123 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
projects/hip/tests/src/hipC.c
-55
Datei anzeigen
@@ -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<stdio.h>
#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<ITER;i++){
Ad[tx] += 1;
}
}
}
int main(){
int A=0, *Ad;
hipMalloc((void**)&Ad, SIZE);
hipMemcpy(Ad, &A, SIZE, hipMemcpyHostToDevice);
dim3 dimGrid, dimBlock;
dimGrid.x = 1, dimGrid.y =1, dimGrid.z = 1;
dimBlock.x = 1, dimBlock.y = 1, dimGrid.z = 1;
hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dimGrid, dimBlock, 0, 0, Ad);
hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
passed();
}
projects/hip/tests/src/hipHcc.cpp
-52
Datei anzeigen
@@ -1,52 +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.
*/
// Test the HCC-specific API extensions for HIP:
/* HIT_START
* BUILD: %t %s EXCLUDE_HIP_PLATFORM nvidia
* TEST: %t EXCLUDE_HIP_PLATFORM all
* HIT_END
*/
#include <stdio.h>
#include <iostream>
#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();
};
projects/hip/tests/src/runtimeApi/event/hipEventMultiThreaded.cpp
+9 -4
Datei anzeigen
@@ -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 <thread>
#include <vector>
#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]);
}
projects/hip/tests/src/runtimeApi/memory/hipManagedKeyword.cpp
+1 -1
Datei anzeigen
@@ -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
projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
+1 -1
Datei anzeigen
@@ -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
*/
projects/hip/tests/src/runtimeApi/module/hipManagedKeyword.cpp
+2 -2
Datei anzeigen
@@ -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
*/
projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp
+1 -1
Datei anzeigen
@@ -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
*/
projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp
+1 -1
Datei anzeigen
@@ -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
*/
projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp
+9 -6
Datei anzeigen
@@ -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 <typename T>
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<int>();
return runTest<int>();
}