AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[dtest] Tests for hipExtModuleLaunchKernel APIs

Browse Source 
Added test scenarios for hipExtModuleLaunchKernel/hipModuleLaunchKernel APIs

 1. Negative Test Scenarios for hipExtModuleLaunchKernel API
 2. Negative Test Scenarios for hipModuleLaunchKernel API
 3. hipExtModuleLaunchKernel API Concurrency Verification
    using global variable
 4. hipExtModuleLaunchKernel API Concurrency Verification
    using total time taken by kernels.
 5. hipExtModuleLaunchKernel API verifying the kernel execution time
    of a particular kernel
 6. hipExtModuleLaunchKernel API verifying by disabling the
    time flag during event creation

SWDEV-238517 for enhancing hip unit tests

Change-Id: I33bbab152d47beba62065c461a07cbcacdc01d97


[ROCm/hip commit: ba3fff4fcd]
This commit is contained in:
DURGESH KROTTAPALLI
2020-11-05 18:44:49 +05:30
committed by Mohan Kumar Mithur
parent facdc585cb
commit d017c25361
3 changed files with 596 additions and 132 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/runtimeApi/module/matmul.cpp
+54 -13
View File
@@ -18,22 +18,63 @@ THE SOFTWARE.
*/
#include"hip/hip_runtime.h"
__device__ int deviceGlobal = 1;
extern "C" __global__ void matmulK(int* A, int* B, int* C, int N) {
int ROW = blockIdx.y*blockDim.y+threadIdx.y;
int COL = blockIdx.x*blockDim.x+threadIdx.x;
int tmpSum = 0;
if (ROW < N && COL < N) {
// each thread computes one element of the block sub-matrix
for (int i = 0; i < N; i++) {
tmpSum += A[ROW * N + i] * B[i * N + COL];
}
int ROW = blockIdx.y*blockDim.y+threadIdx.y;
int COL = blockIdx.x*blockDim.x+threadIdx.x;
int tmpSum = 0;
if ((ROW < N) && (COL < N)) {
// each thread computes one element of the block sub-matrix
for (int i = 0; i < N; i++) {
tmpSum += A[ROW * N + i] * B[i * N + COL];
}
C[ROW * N + COL] =tmpSum;
C[ROW * N + COL] = tmpSum;
}
}
extern "C" __global__ void WaitKernel() {
unsigned long long int wait_t=32000000000,start=clock64(),cur;
do{cur=clock64()-start;}
while(cur<wait_t);
extern "C" __global__ void KernelandExtraParams(int* A, int* B, int* C,
int *D, int N) {
int ROW = blockIdx.y*blockDim.y+threadIdx.y;
int COL = blockIdx.x*blockDim.x+threadIdx.x;
int tmpSum = 0;
if (ROW < N && COL < N) {
// each thread computes one element of the block sub-matrix
for (int i = 0; i < N; i++) {
tmpSum += A[ROW * N + i] * B[i * N + COL];
}
}
C[ROW * N + COL] = tmpSum;
D[ROW * N + COL] = tmpSum;
}
extern "C" __global__ void SixteenSecKernel() {
uint64_t wait_t = 32000000000,
start = clock64(), cur;
do { cur = clock64()-start;}while (cur < wait_t);
}
extern "C" __global__ void TwoSecKernel() {
if (deviceGlobal == 0x2222) {
deviceGlobal = 0x3333;
}
uint64_t wait_t = 4000000000,
start = clock64(), cur;
do { cur = clock64()-start;}while (cur < wait_t);
if (deviceGlobal != 0x3333) {
deviceGlobal = 0x5555;
}
}
extern "C" __global__ void FourSecKernel() {
if (deviceGlobal == 1) {
deviceGlobal = 0x2222;
}
uint64_t wait_t = 8000000000,
start = clock64(), cur;
do { cur = clock64()-start;}while (cur < wait_t);
if (deviceGlobal == 0x2222) {
deviceGlobal = 0x4444;
}
}