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

SWDEV-337452 - Changing Clock64 to WallClock64 in directed tests. (#3140)

Browse Source 
Change-Id: I511ab4dcc61daee4fdfbd2a248b5fe74e52174b2


[ROCm/hip commit: 0ea181501c]
This commit is contained in:
ROCm CI Service Account
2023-03-21 01:48:32 +05:30
committed by GitHub
parent dd539e475d
commit daf4de3aa4
18 changed files with 604 additions and 58 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp
+21 -1
View File
@@ -61,6 +61,25 @@ __global__ void vector_square(float* C_d, float* A_d, size_t N_elmts) {
}
}
__global__ void vector_square_gfx11(float* C_d, float* A_d, size_t N_elmts) {
#ifdef __HIP_PLATFORM_AMD__
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = offset; i < N_elmts; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
// Delay the thread 1
if (offset == 1) {
unsigned long long int wait_t = 3200000000, start = wall_clock64(), cur;
do {
cur = wall_clock64() - start;
} while (cur < wait_t);
}
#endif
}
float *A_h, *C_h;
static void HIPRT_CB Callback1(hipStream_t stream, hipError_t status,
@@ -102,7 +121,8 @@ int main(int argc, char* argv[]) {
const unsigned threadsPerBlock = 256;
const unsigned blocks = (N_elmts + 255)/threadsPerBlock;
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0,
auto vector_square_used = IsGfx11() ? vector_square_gfx11 : vector_square;
hipLaunchKernelGGL((vector_square_used), dim3(blocks), dim3(threadsPerBlock), 0,
mystream, C_d, A_d, N_elmts);
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
HIPCHECK(hipStreamAddCallback(mystream, Callback1, NULL, 0));