summarylogtreecommitdiffstats
diff options
context:
space:
mode:
authorTorsten Keßler2022-12-07 10:03:41 +0100
committerTorsten Keßler2022-12-07 10:03:41 +0100
commit0875c7cd005ef6696ed8271e612e78d905f011f0 (patch)
tree13afa8b7eafd787d40b2a582782e7f389ba92eae
parent8557edb951ed3609af6b6b229af3d6bc44a5e3ed (diff)
downloadaur-rocm-hip-runtime.tar.gz
[hip-runtime-amd] Add simple test case
-rw-r--r--test.cpp66
-rwxr-xr-xtest.sh6
2 files changed, 72 insertions, 0 deletions
diff --git a/test.cpp b/test.cpp
new file mode 100644
index 000000000000..1f2e903e1a23
--- /dev/null
+++ b/test.cpp
@@ -0,0 +1,66 @@
+#include <iostream>
+#include <cmath>
+#include <vector>
+#include <hip/hip_runtime.h>
+
+__global__
+void saxpy(int n, float a, const float *x, float *y)
+{
+ int i = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+ if(i < n){
+ y[i] = a * x[i] + y[i];
+ }
+}
+
+__global__
+void sset(int n, float a, float *x)
+{
+ int i = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+ if(i < n){
+ x[i] = a;
+ }
+}
+
+int main()
+{
+ hipDeviceProp_t prop;
+ hipGetDeviceProperties(&prop, 0);
+ std::cout << "Agent " << prop.name << "\n";
+ std::cout << "System version " << prop.major
+ << "." << prop.minor << "\n";
+
+ int n = 1024;
+ float *x;
+ float *y;
+ hipMalloc((void**)&x, sizeof *x * n);
+ hipMalloc((void**)&y, sizeof *y * n);
+
+ std::vector<float> xin(n);
+ for(int i = 0; i < n; i++){
+ xin[i] = -1.0 + 2.0 * i / n;
+ }
+ hipMemcpy(x, xin.data(), sizeof *x * n, hipMemcpyHostToDevice);
+
+ float ac = -14.412f;
+ hipLaunchKernelGGL(sset, dim3(1), dim3(n), 0, 0, n, ac, y);
+
+ float a = 5321.124f;
+ hipLaunchKernelGGL(saxpy, dim3(1), dim3(n), 0, 0, n, a, x, y);
+
+ std::vector<float> yout(n);
+ hipMemcpy(yout.data(), y, sizeof *y * n, hipMemcpyDeviceToHost);
+
+ hipFree(x);
+ hipFree(y);
+
+ for(int i = 0; i < n; i++){
+ yout[i] -= a * xin[i];
+ if(std::abs(yout[i] - ac) > 0.001f){
+ std::cout << "Test failed at index " << i
+ << " with entry " << yout[i]
+ << " (" << ac << ")\n";
+ return 1;
+ }
+ }
+ std::cout << "TESTS PASSED!" << std::endl;
+}
diff --git a/test.sh b/test.sh
new file mode 100755
index 000000000000..8f17dd769682
--- /dev/null
+++ b/test.sh
@@ -0,0 +1,6 @@
+#!/usr/bin/env sh
+
+OUT=$(mktemp -d)
+
+/opt/rocm/bin/hipcc -o "$OUT/test" test.cpp
+"$OUT"/test