Decoupling HIP code with runFeatureTest target

Change-Id: I8cac692289c7c9393b76789b0fa67eb239df2079

tests-v2/featuretests/profiler/apps/hip_kernels.cpp

@@ -0,0 +1,106 @@
+#include <hip/hip_runtime.h>
+#ifdef NDEBUG
+#define HIP_ASSERT(x) x
+#else
+#define HIP_ASSERT(x) (assert((x) == hipSuccess))
+#endif
+
+#define WIDTH 1024
+#define HEIGHT 1024
+
+#define NUM (WIDTH * HEIGHT)
+
+#define THREADS_PER_BLOCK_X 16
+#define THREADS_PER_BLOCK_Y 16
+#define THREADS_PER_BLOCK_Z 1
+
+// empty kernel
+__global__ void kernel() {}
+
+// vector add kernel
+__global__ void vectoradd_float(float* __restrict__ a, const float* __restrict__ b,
+                                const float* __restrict__ c, int width, int height)
+
+{
+  int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+  int i = y * width + x;
+  if (i < (width * height)) {
+    a[i] = b[i] + c[i];
+  }
+}
+
+// launches an empty kernel in profiler context
+void KernelLaunch() {
+  // run empty kernel
+  kernel<<<1, 1>>>();
+  hipDeviceSynchronize();
+}
+
+int LaunchVectorAddKernel() {
+  float* hostA;
+  float* hostB;
+  float* hostC;
+
+  float* deviceA;
+  float* deviceB;
+  float* deviceC;
+
+  hipDeviceProp_t devProp;
+  hipGetDeviceProperties(&devProp, 0);
+
+  int i;
+  int errors;
+
+  hostA = (float*)malloc(NUM * sizeof(float));
+  hostB = (float*)malloc(NUM * sizeof(float));
+  hostC = (float*)malloc(NUM * sizeof(float));
+
+  // initialize the input data
+  for (i = 0; i < NUM; i++) {
+    hostB[i] = (float)i;
+    hostC[i] = (float)i * 100.0f;
+  }
+
+  HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float)));
+  HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float)));
+  HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float)));
+
+  HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice));
+  HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice));
+
+
+  hipLaunchKernelGGL(vectoradd_float,
+                     dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
+                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
+                     deviceC, WIDTH, HEIGHT);
+
+
+  HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostA[i] != (hostB[i] + hostC[i])) {
+      errors++;
+    }
+  }
+  if (errors != 0) {
+    printf("FAILED: %d errors\n", errors);
+  } else {
+    printf("PASSED!\n");
+  }
+
+  HIP_ASSERT(hipFree(deviceA));
+  HIP_ASSERT(hipFree(deviceB));
+  HIP_ASSERT(hipFree(deviceC));
+
+  free(hostA);
+  free(hostB);
+  free(hostC);
+
+  // hipResetDefaultAccelerator();
+
+  return errors;
+}