Add more tests

Olympus-HPC · Nov 26, 2024 · 70d10d6 · 70d10d6
1 parent 8557bde
commit 70d10d6
Show file tree

Hide file tree

Showing 4 changed files with 143 additions and 0 deletions.
diff --git a/tests/gpu/CMakeLists.txt b/tests/gpu/CMakeLists.txt
@@ -114,8 +114,11 @@ CREATE_GPU_TEST(kernels_gvar kernels_gvar.cpp)
 CREATE_GPU_TEST(kernel_launches kernel_launches.cpp)
 CREATE_GPU_TEST(kernel_launches_args kernel_launches_args.cpp)
 CREATE_GPU_TEST(indirect_launcher indirect_launcher.cpp)
+CREATE_GPU_TEST(indirect_launcher_arg indirect_launcher_arg.cpp)
 CREATE_GPU_TEST(indirect_launcher_tpl_multi indirect_launcher_tpl_multi.cpp)
+CREATE_GPU_TEST(indirect_launcher_tpl_multi_arg indirect_launcher_tpl_multi_arg.cpp)
 CREATE_GPU_TEST(indirect_launcher_multi indirect_launcher_multi.cpp)
+CREATE_GPU_TEST(indirect_launcher_multi_arg indirect_launcher_multi_arg.cpp)
 CREATE_GPU_TEST(multi_file file1_kernel.cpp file2_kernel.cpp)
 CREATE_GPU_TEST(daxpy daxpy.cpp)
 CREATE_GPU_TEST(kernel_host_jit kernel_host_jit.cpp)

diff --git a/tests/gpu/indirect_launcher_arg.cpp b/tests/gpu/indirect_launcher_arg.cpp
@@ -0,0 +1,43 @@
+// clang-format off
+// RUN: ./indirect_launcher.%ext | FileCheck %s --check-prefixes=CHECK,CHECK-FIRST
+// Second run uses the object cache.
+// RUN: ./indirect_launcher.%ext | FileCheck %s --check-prefixes=CHECK,CHECK-SECOND
+// clang-format on
+
+#include <climits>
+#include <cstdio>
+
+#include "gpu_common.h"
+
+#define gpuErrCheck(CALL)                                                      \
+  {                                                                            \
+    gpuError_t err = CALL;                                                     \
+    if (err != gpuSuccess) {                                                   \
+      printf("ERROR @ %s:%d ->  %s\n", __FILE__, __LINE__,                     \
+             gpuGetErrorString(err));                                          \
+      abort();                                                                 \
+    }                                                                          \
+  }
+
+__global__ __attribute__((annotate("jit", 1))) void kernel(int a) {
+  printf("Kernel %d\n", a);
+}
+
+template <typename T> gpuError_t launcher(T kernel_in, int a) {
+  void *args[] = {&a};
+  return gpuLaunchKernel((const void *)kernel_in, 1, 1, args, 0, 0);
+}
+
+int main() {
+  kernel<<<1, 1>>>(42);
+  gpuErrCheck(launcher(kernel, 24));
+  gpuErrCheck(gpuDeviceSynchronize());
+  return 0;
+}
+
+// CHECK: Kernel 42
+// CHECK: Kernel 24
+// CHECK: JitCache hits 0 total 2
+// CHECK: HashValue {{[0-9]+}} NumExecs 1 NumHits 0
+// CHECK-FIRST: JitStorageCache hits 0 total 1
+// CHECK-SECOND: JitStorageCache hits 1 total 1
diff --git a/tests/gpu/indirect_launcher_multi_arg.cpp b/tests/gpu/indirect_launcher_multi_arg.cpp
@@ -0,0 +1,47 @@
+// clang-format off
+// RUN: ./indirect_launcher.%ext | FileCheck %s --check-prefixes=CHECK,CHECK-FIRST
+// Second run uses the object cache.
+// RUN: ./indirect_launcher.%ext | FileCheck %s --check-prefixes=CHECK,CHECK-SECOND
+// clang-format on
+
+#include <climits>
+#include <cstdio>
+
+#include "gpu_common.h"
+
+#define gpuErrCheck(CALL)                                                      \
+  {                                                                            \
+    gpuError_t err = CALL;                                                     \
+    if (err != gpuSuccess) {                                                   \
+      printf("ERROR @ %s:%d ->  %s\n", __FILE__, __LINE__,                     \
+             gpuGetErrorString(err));                                          \
+      abort();                                                                 \
+    }                                                                          \
+  }
+
+__global__ __attribute__((annotate("jit", 1)))void kernel(int arg) {
+  printf("Kernel one; arg = %d\n", arg);
+}
+
+__global__ __attribute__((annotate("jit", 1))) void kernel_two(int arg) {
+  printf("Kernel two; arg = %d\n", arg);
+}
+
+gpuError_t launcher(void** kernel_in, int a) {
+  void *args[] = {&a};
+  return gpuLaunchKernel((const void *)kernel_in, 1, 1, args, 0, 0);
+}
+
+int main() {
+  gpuErrCheck(launcher((void**)kernel, 42));
+  gpuErrCheck(launcher((void**)kernel_two, 24));
+  gpuErrCheck(gpuDeviceSynchronize());
+  return 0;
+}
+
+// CHECK: Kernel one; arg = 42
+// CHECK: Kernel two; arg = 24
+// CHECK: JitCache hits 0 total 2
+// CHECK: HashValue {{[0-9]+}} NumExecs 1 NumHits 0
+// CHECK-FIRST: JitStorageCache hits 0 total 2
+// CHECK-SECOND: JitStorageCache hits 2 total 2
diff --git a/tests/gpu/indirect_launcher_tpl_multi_arg.cpp b/tests/gpu/indirect_launcher_tpl_multi_arg.cpp
@@ -0,0 +1,50 @@
+// clang-format off
+// RUN: ./indirect_launcher.%ext | FileCheck %s --check-prefixes=CHECK,CHECK-FIRST
+// Second run uses the object cache.
+// RUN: ./indirect_launcher.%ext | FileCheck %s --check-prefixes=CHECK,CHECK-SECOND
+// clang-format on
+
+#include <climits>
+#include <cstdio>
+#include <iostream>
+
+#include "gpu_common.h"
+
+#define gpuErrCheck(CALL)                                                      \
+  {                                                                            \
+    gpuError_t err = CALL;                                                     \
+    if (err != gpuSuccess) {                                                   \
+      printf("ERROR @ %s:%d ->  %s\n", __FILE__, __LINE__,                     \
+             gpuGetErrorString(err));                                          \
+      abort();                                                                 \
+    }                                                                          \
+  }
+
+__global__ __attribute__((annotate("jit", 1)))void kernel(int arg) {
+  printf("Kernel one; arg = %d\n", arg);
+}
+
+__global__ __attribute__((annotate("jit", 1))) void kernel_two(int arg) {
+  printf("Kernel two; arg = %d\n", arg);
+}
+
+template <typename T> gpuError_t launcher(T kernel_in, int a) {
+  void *args[] = {&a};
+  return gpuLaunchKernel((const void*)kernel_in, 1, 1, args, 0, 0);
+}
+
+int main() {
+  auto indirect = reinterpret_cast<const void*>(&kernel);
+  gpuErrCheck(launcher(kernel, 42));
+  gpuErrCheck(launcher(kernel_two, 24));
+  gpuErrCheck(gpuDeviceSynchronize());
+
+  return 0;
+}
+
+// CHECK: Kernel one; arg = 42
+// CHECK: Kernel two; arg = 24
+// CHECK: JitCache hits 0 total 2
+// CHECK: HashValue {{[0-9]+}} NumExecs 1 NumHits 0
+// CHECK-FIRST: JitStorageCache hits 0 total 2
+// CHECK-SECOND: JitStorageCache hits 2 total 2