ROCmWrapper/plugins/ROCmTestWrapperAdditionModule.cc

0001 #include <cstddef>
0002 #include <cstdint>
0003 #include <iostream>
0004 #include <random>
0005 #include <vector>
0006
0007 #include <hip/hip_runtime.h>
0008
0009 #include "FWCore/Framework/interface/Event.h"
0010 #include "FWCore/Framework/interface/Frameworkfwd.h"
0011 #include "FWCore/Framework/interface/global/EDAnalyzer.h"
0012 #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
0013 #include "FWCore/ParameterSet/interface/ParameterSet.h"
0014 #include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
0015 #include "FWCore/ServiceRegistry/interface/Service.h"
0016 #include "HeterogeneousCore/ROCmServices/interface/ROCmInterface.h"
0017 #include "HeterogeneousTest/ROCmWrapper/interface/DeviceAdditionWrapper.h"
0018 #include "HeterogeneousCore/ROCmUtilities/interface/hipCheck.h"
0019
0020 class ROCmTestWrapperAdditionModule : public edm::global::EDAnalyzer<> {
0021 public:
0022   explicit ROCmTestWrapperAdditionModule(edm::ParameterSet const& config);
0023   ~ROCmTestWrapperAdditionModule() override = default;
0024
0025   static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
0026
0027   void analyze(edm::StreamID, edm::Event const& event, edm::EventSetup const& setup) const override;
0028
0029 private:
0030   const uint32_t size_;
0031 };
0032
0033 ROCmTestWrapperAdditionModule::ROCmTestWrapperAdditionModule(edm::ParameterSet const& config)
0034     : size_(config.getParameter<uint32_t>("size")) {}
0035
0036 void ROCmTestWrapperAdditionModule::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
0037   edm::ParameterSetDescription desc;
0038   desc.add<uint32_t>("size", 1024 * 1024);
0039   descriptions.addWithDefaultLabel(desc);
0040 }
0041
0042 void ROCmTestWrapperAdditionModule::analyze(edm::StreamID,
0043                                             edm::Event const& event,
0044                                             edm::EventSetup const& setup) const {
0045   // require ROCm for running
0046   edm::Service<ROCmInterface> rocm;
0047   if (not rocm or not rocm->enabled()) {
0048     std::cout << "The ROCmService is not available or disabled, the test will be skipped.\n";
0049     return;
0050   }
0051
0052   // random number generator with a gaussian distribution
0053   std::random_device rd{};
0054   std::default_random_engine rand{rd()};
0055   std::normal_distribution<float> dist{0., 1.};
0056
0057   // tolerance
0058   constexpr float epsilon = 0.000001;
0059
0060   // allocate input and output host buffers
0061   std::vector<float> in1_h(size_);
0062   std::vector<float> in2_h(size_);
0063   std::vector<float> out_h(size_);
0064
0065   // fill the input buffers with random data, and the output buffer with zeros
0066   for (size_t i = 0; i < size_; ++i) {
0067     in1_h[i] = dist(rand);
0068     in2_h[i] = dist(rand);
0069     out_h[i] = 0.;
0070   }
0071
0072   // allocate input and output buffers on the device
0073   float* in1_d;
0074   float* in2_d;
0075   float* out_d;
0076   hipCheck(hipMalloc(&in1_d, size_ * sizeof(float)));
0077   hipCheck(hipMalloc(&in2_d, size_ * sizeof(float)));
0078   hipCheck(hipMalloc(&out_d, size_ * sizeof(float)));
0079
0080   // copy the input data to the device
0081   hipCheck(hipMemcpy(in1_d, in1_h.data(), size_ * sizeof(float), hipMemcpyHostToDevice));
0082   hipCheck(hipMemcpy(in2_d, in2_h.data(), size_ * sizeof(float), hipMemcpyHostToDevice));
0083
0084   // fill the output buffer with zeros
0085   hipCheck(hipMemset(out_d, 0, size_ * sizeof(float)));
0086
0087   // launch the 1-dimensional kernel for vector addition
0088   cms::rocmtest::wrapper_add_vectors_f(in1_d, in2_d, out_d, size_);
0089
0090   // copy the results from the device to the host
0091   hipCheck(hipMemcpy(out_h.data(), out_d, size_ * sizeof(float), hipMemcpyDeviceToHost));
0092
0093   // wait for all the operations to complete
0094   hipCheck(hipDeviceSynchronize());
0095
0096   // check the results
0097   for (size_t i = 0; i < size_; ++i) {
0098     float sum = in1_h[i] + in2_h[i];
0099     assert(out_h[i] < sum + epsilon);
0100     assert(out_h[i] > sum - epsilon);
0101   }
0102
0103   std::cout << "All tests passed.\n";
0104 }
0105
0106 #include "FWCore/Framework/interface/MakerMacros.h"
0107 DEFINE_FWK_MODULE(ROCmTestWrapperAdditionModule);