Add indirect access memory stress test

Signed-off-by: Oskar Hubert Weber <[email protected]>

Author: oskarweb

stress_tests/common/src/stress_common_func.cpp

@@ -1,6 +1,6 @@
 /*
  *
- * Copyright (C) 2019-2023 Intel Corporation
+ * Copyright (C) 2019-2025 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -43,7 +43,7 @@ uint64_t total_available_host_memory() {
   stat.dwLength = sizeof(stat);
   GlobalMemoryStatusEx(&stat);
 
-  return stat.ullAvailVirtual;
+  return stat.ullAvailPhys;
 }
 uint64_t get_page_size() {
   SYSTEM_INFO si;

stress_tests/test_memory_allocation/kernels/test_multiple_memory_allocations.cl

@@ -1,6 +1,6 @@
 /*
  *
- * Copyright (C) 2019 Intel Corporation
+ * Copyright (C) 2019-2025 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -86,3 +86,78 @@ kernel void test_device_memory10_unit_size4(__global uint *src, __global uint *d
     size_t  tid = get_global_id(0);
     dst[tid] = src[tid];
 }
+
+struct buffer {
+  uint *data;
+};
+
+kernel void test_device_memory1_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory2_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory3_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory4_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory5_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory6_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory7_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory8_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory9_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+
+kernel void test_device_memory10_indirect(__global struct buffer *src_ptrs, __global struct buffer *dst_ptrs, uint dispatch_id) {
+	uint *src = src_ptrs[dispatch_id].data;
+    uint *dst = dst_ptrs[dispatch_id].data;
+    size_t tid = get_global_id(0);
+    dst[tid] = src[tid];
+}
+