AzIntEngineGPU: Fixes
This commit is contained in:
@@ -100,13 +100,13 @@ AzIntEngineGPU::AzIntEngineGPU(const AzimuthalIntegration &integration)
|
||||
cpu_count_reg(azint_count),
|
||||
gpu_image(npixel) {
|
||||
|
||||
cudaDeviceProp prop;
|
||||
cudaDeviceProp prop{};
|
||||
cudaGetDeviceProperties(&prop, 0);
|
||||
|
||||
threads = 128;
|
||||
blocks = 4 * prop.multiProcessorCount;
|
||||
shared_size = prop.sharedMemPerBlock;
|
||||
shared_needed = azint_bins * sizeof(float) + azint_bins * sizeof(float) + azint_bins * sizeof(uint32_t);
|
||||
shared_needed = azint_bins * (2 * sizeof(float) + sizeof(uint32_t));
|
||||
|
||||
cudaMemcpy(gpu_azint_correction, integration.Corrections().data(), sizeof(float) * npixel,
|
||||
cudaMemcpyHostToDevice);
|
||||
@@ -124,25 +124,13 @@ void AzIntEngineGPU::Run(const std::vector<int32_t> &image, AzimuthalIntegration
|
||||
|
||||
if (shared_needed < shared_size) {
|
||||
gpu_azim_shared<<<blocks, threads, shared_needed, stream>>>(
|
||||
gpu_pixel_to_bin,
|
||||
gpu_azint_correction,
|
||||
gpu_image,
|
||||
gpu_sum,
|
||||
gpu_sum2,
|
||||
gpu_count,
|
||||
npixel,
|
||||
azint_bins
|
||||
gpu_pixel_to_bin,gpu_azint_correction,gpu_image, gpu_sum, gpu_sum2,
|
||||
gpu_count, npixel, azint_bins
|
||||
);
|
||||
} else {
|
||||
gpu_azim<<<blocks, threads, 0, stream>>>(
|
||||
gpu_pixel_to_bin,
|
||||
gpu_azint_correction,
|
||||
gpu_image,
|
||||
gpu_sum,
|
||||
gpu_sum2,
|
||||
gpu_count,
|
||||
npixel,
|
||||
azint_bins
|
||||
gpu_pixel_to_bin,gpu_azint_correction,gpu_image, gpu_sum, gpu_sum2,
|
||||
gpu_count, npixel, azint_bins
|
||||
);
|
||||
}
|
||||
|
||||
|
||||
@@ -10,8 +10,8 @@ class AzIntEngineGPU : public AzIntEngine {
|
||||
CudaStream stream;
|
||||
int threads;
|
||||
int blocks;
|
||||
int shared_needed;
|
||||
int shared_size;
|
||||
size_t shared_needed;
|
||||
size_t shared_size;
|
||||
|
||||
CudaDevicePtr<float> gpu_azint_correction;
|
||||
CudaDevicePtr<uint16_t> gpu_pixel_to_bin;
|
||||
|
||||
Reference in New Issue
Block a user