ROCmKernel/test/testDeviceAdditionKernel.hip.cc

0001 #include <cstddef>
0002 #include <cstdint>
0003 #include <random>
0004 #include <vector>
0005
0006 #define CATCH_CONFIG_MAIN
0007 #include <catch.hpp>
0008
0009 #include <hip/hip_runtime.h>
0010
0011 #include "HeterogeneousTest/ROCmKernel/interface/DeviceAdditionKernel.h"
0012 #include "HeterogeneousCore/ROCmUtilities/interface/hipCheck.h"
0013 #include "HeterogeneousCore/ROCmUtilities/interface/requireDevices.h"
0014
0015 TEST_CASE("HeterogeneousTest/ROCmKernel test", "[rocmTestKernelAdditionKernel]") {
0016   cms::rocmtest::requireDevices();
0017
0018   // random number generator with a gaussian distribution
0019   std::random_device rd{};
0020   std::default_random_engine rand{rd()};
0021   std::normal_distribution<float> dist{0., 1.};
0022
0023   // tolerance
0024   constexpr float epsilon = 0.000001;
0025
0026   // buffer size
0027   constexpr size_t size = 1024 * 1024;
0028
0029   // allocate input and output host buffers
0030   std::vector<float> in1_h(size);
0031   std::vector<float> in2_h(size);
0032   std::vector<float> out_h(size);
0033
0034   // fill the input buffers with random data, and the output buffer with zeros
0035   for (size_t i = 0; i < size; ++i) {
0036     in1_h[i] = dist(rand);
0037     in2_h[i] = dist(rand);
0038     out_h[i] = 0.;
0039   }
0040
0041   SECTION("Test add_vectors_f") {
0042     // allocate input and output buffers on the device
0043     float* in1_d;
0044     float* in2_d;
0045     float* out_d;
0046     REQUIRE_NOTHROW(hipCheck(hipMalloc(&in1_d, size * sizeof(float))));
0047     REQUIRE_NOTHROW(hipCheck(hipMalloc(&in2_d, size * sizeof(float))));
0048     REQUIRE_NOTHROW(hipCheck(hipMalloc(&out_d, size * sizeof(float))));
0049
0050     // copy the input data to the device
0051     REQUIRE_NOTHROW(hipCheck(hipMemcpy(in1_d, in1_h.data(), size * sizeof(float), hipMemcpyHostToDevice)));
0052     REQUIRE_NOTHROW(hipCheck(hipMemcpy(in2_d, in2_h.data(), size * sizeof(float), hipMemcpyHostToDevice)));
0053
0054     // fill the output buffer with zeros
0055     REQUIRE_NOTHROW(hipCheck(hipMemset(out_d, 0, size * sizeof(float))));
0056
0057     // launch the 1-dimensional kernel for vector addition
0058     cms::rocmtest::kernel_add_vectors_f<<<32, 32>>>(in1_d, in2_d, out_d, size);
0059     REQUIRE_NOTHROW(hipCheck(hipGetLastError()));
0060
0061     // copy the results from the device to the host
0062     REQUIRE_NOTHROW(hipCheck(hipMemcpy(out_h.data(), out_d, size * sizeof(float), hipMemcpyDeviceToHost)));
0063
0064     // wait for all the operations to complete
0065     REQUIRE_NOTHROW(hipCheck(hipDeviceSynchronize()));
0066
0067     // check the results
0068     for (size_t i = 0; i < size; ++i) {
0069       float sum = in1_h[i] + in2_h[i];
0070       CHECK_THAT(out_h[i], Catch::Matchers::WithinAbs(sum, epsilon));
0071     }
0072   }
0073 }