diff --git a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPU.cc b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPU.cc index e38c596fbe2f5..8fe5688018728 100644 --- a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPU.cc +++ b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPU.cc @@ -29,7 +29,8 @@ class TestCUDAAnalyzerGPU : public edm::global::EDAnalyzer<> { edm::EDGetTokenT> const srcToken_; double const minValue_; double const maxValue_; - std::unique_ptr gpuAlgo_; + // the public interface is thread safe + CMS_THREAD_SAFE mutable std::unique_ptr gpuAlgo_; }; TestCUDAAnalyzerGPU::TestCUDAAnalyzerGPU(edm::ParameterSet const& iConfig) diff --git a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.cu b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.cu index 4d4cca09e4668..01ded40c6d7ff 100644 --- 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<<>>(d_input, sum_.get(), NUM_VALUES); } diff --git a/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.h b/HeterogeneousCore/CUDATest/plugins/TestCUDAAnalyzerGPUKernel.h index 6854ba8d61af7..612e617c67c8c 100644 --- 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 sum_; // all writes are atomic in CUDA + cudautils::device::unique_ptr sum_; // all writes are atomic in CUDA }; #endif