Move mutability of a member from TestCUDAAnalyzerGPUKernel to TestCUD…

…AAnalyzerGPU The point is that the member functions mutating the visible state should not be const even if they are thread safe. TestCUDAAnalyzerGPUKernel::analyzeAsync() mutates the visible state (by "filling a histogram"), so it should not be const. Declaring TestCUDAAnalyzerGPU::gpuAlgo_ as mutable is an improvement since filling the histogram does not really change the visible state of the EDAnalyzer (towards the framework).
cms-patatrack · Dec 13, 2019 · 0a5c1ab · 0a5c1ab
1 parent d6b4490
commit 0a5c1ab
Show file tree

Hide file tree

Showing 3 changed files with 6 additions and 5 deletions.
diff --git a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPU.cc b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPU.cc
@@ -29,7 +29,8 @@ class TestCUDAAnalyzerGPU : public edm::global::EDAnalyzer<> {
   edm::EDGetTokenT<CUDAProduct<CUDAThing>> const srcToken_;
   double const minValue_;
   double const maxValue_;
-  std::unique_ptr<TestCUDAAnalyzerGPUKernel> gpuAlgo_;
+  // the public interface is thread safe
+  CMS_THREAD_SAFE mutable std::unique_ptr<TestCUDAAnalyzerGPUKernel> gpuAlgo_;
 };
 
 TestCUDAAnalyzerGPU::TestCUDAAnalyzerGPU(edm::ParameterSet const& iConfig)

diff --git a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.cu b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.cu
@@ -31,7 +31,7 @@ TestCUDAAnalyzerGPUKernel::TestCUDAAnalyzerGPUKernel(cudaStream_t stream) {
   cudaCheck(cudaStreamSynchronize(stream));
 }
 
-void TestCUDAAnalyzerGPUKernel::analyzeAsync(const float *d_input, cudaStream_t stream) const {
+void TestCUDAAnalyzerGPUKernel::analyzeAsync(const float *d_input, cudaStream_t stream) {
   analyze<<<int(ceil(float(NUM_VALUES) / 256)), 256, 0, stream>>>(d_input, sum_.get(), NUM_VALUES);
 }
 

diff --git a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.h b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.h
@@ -12,12 +12,12 @@ class TestCUDAAnalyzerGPUKernel {
   TestCUDAAnalyzerGPUKernel(cudaStream_t stream);
   ~TestCUDAAnalyzerGPUKernel() = default;
 
-  // returns (owning) pointer to device memory
-  void analyzeAsync(const float* d_input, cudaStream_t stream) const;
+  // thread safe
+  void analyzeAsync(const float* d_input, cudaStream_t stream);
   float value(cudaStream_t stream) const;
 
 private:
-  mutable cudautils::device::unique_ptr<float[]> sum_;  // all writes are atomic in CUDA
+  cudautils::device::unique_ptr<float[]> sum_;  // all writes are atomic in CUDA
 };
 
 #endif