Skip to content

Commit 420c622

Browse files
davidrohrdavidrohr
committed
GPU: Replace double constants by float constants
1 parent 023cef5 commit 420c622

46 files changed

Lines changed: 161 additions & 161 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎GPU/Common/test/testGPUsortCUDA.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ struct TestEnvironment {
4444

4545
// create an array of unordered floats with negative and positive values
4646
for (size_t i = 0; i < size; i++) {
47-
data[i] = size / 2.0 - i;
47+
data[i] = size / 2.0f - i;
4848
}
4949
// create copy
5050
std::memcpy(sorted.data(), data, size * sizeof(float));

‎GPU/GPUTracking/Base/GPUParam.cxx‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,7 @@ void GPUParam::SetDefaults(float solenoidBz)
9191
par.dAlpha = 0.349066f;
9292
bzkG = solenoidBz;
9393
constBz = bzkG * GPUCA_NAMESPACE::gpu::gpu_common_constants::kCLight;
94-
qptB5Scaler = CAMath::Abs(bzkG) > 0.1 ? CAMath::Abs(bzkG) / 5.006680f : 1.f;
94+
qptB5Scaler = CAMath::Abs(bzkG) > 0.1f ? CAMath::Abs(bzkG) / 5.006680f : 1.f;
9595
par.dodEdx = 0;
9696

9797
constexpr float plusZmin = 0.0529937;
@@ -109,7 +109,7 @@ void GPUParam::SetDefaults(float solenoidBz)
109109
if (tmp >= GPUCA_NSLICES / 4) {
110110
tmp -= GPUCA_NSLICES / 2;
111111
}
112-
SliceParam[i].Alpha = 0.174533 + par.dAlpha * tmp;
112+
SliceParam[i].Alpha = 0.174533f + par.dAlpha * tmp;
113113
SliceParam[i].CosAlpha = CAMath::Cos(SliceParam[i].Alpha);
114114
SliceParam[i].SinAlpha = CAMath::Sin(SliceParam[i].Alpha);
115115
SliceParam[i].AngleMin = SliceParam[i].Alpha - par.dAlpha / 2.f;
@@ -145,7 +145,7 @@ void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessi
145145
}
146146
}
147147
par.earlyTpcTransform = rec.tpc.forceEarlyTransform == -1 ? (!par.continuousTracking) : rec.tpc.forceEarlyTransform;
148-
qptB5Scaler = CAMath::Abs(bzkG) > 0.1 ? CAMath::Abs(bzkG) / 5.006680f : 1.f;
148+
qptB5Scaler = CAMath::Abs(bzkG) > 0.1f ? CAMath::Abs(bzkG) / 5.006680f : 1.f;
149149
if (p) {
150150
par.debugLevel = p->debugLevel;
151151
par.resetTimers = p->resetTimers;

‎GPU/GPUTracking/Base/GPUReconstructionTimeframe.cxx‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ extern GPUSettingsStandalone configStandalone;
3939
}
4040
static auto& config = configStandalone.TF;
4141

42-
GPUReconstructionTimeframe::GPUReconstructionTimeframe(GPUChainTracking* chain, int (*read)(int), int nEvents) : mChain(chain), mReadEvent(read), mNEventsInDirectory(nEvents), mDisUniReal(0., 1.), mRndGen1(configStandalone.seed), mRndGen2(mDisUniInt(mRndGen1))
42+
GPUReconstructionTimeframe::GPUReconstructionTimeframe(GPUChainTracking* chain, int (*read)(int), int nEvents) : mChain(chain), mReadEvent(read), mNEventsInDirectory(nEvents), mDisUniReal(0.f, 1.f), mRndGen1(configStandalone.seed), mRndGen2(mDisUniInt(mRndGen1))
4343
{
4444
mMaxBunchesFull = TIME_ORBIT / config.bunchSpacing;
4545
mMaxBunches = (TIME_ORBIT - config.abortGapTime) / config.bunchSpacing;
@@ -79,7 +79,7 @@ int GPUReconstructionTimeframe::ReadEventShifted(int iEvent, float shiftZ, float
7979
}
8080
}
8181
}
82-
if (shiftZ != 0.) {
82+
if (shiftZ != 0.f) {
8383
for (unsigned int iSlice = 0; iSlice < NSLICES; iSlice++) {
8484
for (unsigned int j = 0; j < mChain->mIOPtrs.nClusterData[iSlice]; j++) {
8585
auto& tmp = mChain->mIOMem.clusterData[iSlice][j];
@@ -339,18 +339,18 @@ int GPUReconstructionTimeframe::LoadMergedEvents(int iEvent)
339339
if (iEventInTimeframe == 0) {
340340
shift = 0;
341341
} else {
342-
shift = (iEventInTimeframe - 0.5 + shift) * config.averageDistance;
342+
shift = (iEventInTimeframe - 0.5f + shift) * config.averageDistance;
343343
}
344344
}
345345
} else {
346346
if (config.shiftFirstEvent) {
347-
shift = config.averageDistance * (iEventInTimeframe + 0.5);
347+
shift = config.averageDistance * (iEventInTimeframe + 0.5f);
348348
} else {
349349
shift = config.averageDistance * (iEventInTimeframe);
350350
}
351351
}
352352
} else {
353-
shift = 0.;
353+
shift = 0.f;
354354
}
355355

356356
if (ReadEventShifted(iEvent * config.nMerge + iEventInTimeframe, shift) < 0) {

‎GPU/GPUTracking/Base/GPUReconstructionTimeframe.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ class GPUReconstructionTimeframe
6262
std::mt19937_64 mRndGen2;
6363

6464
int mTrainDist = 0;
65-
float mCollisionProbability = 0.;
65+
float mCollisionProbability = 0.f;
6666
int mMaxBunchesFull;
6767
int mMaxBunches;
6868

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAInternals.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ class GPUDebugTiming
6969
mRec->GPUFailedMsg(cudaEventSynchronize((cudaEvent_t)mDeviceTimers[1]));
7070
float v;
7171
mRec->GPUFailedMsg(cudaEventElapsedTime(&v, (cudaEvent_t)mDeviceTimers[0], (cudaEvent_t)mDeviceTimers[1]));
72-
mXYZ.t = v * 1.e-3;
72+
mXYZ.t = v * 1.e-3f;
7373
} else {
7474
mRec->GPUFailedMsg(cudaStreamSynchronize(mStreams[mXYZ.x.stream]));
7575
mXYZ.t = mTimer.GetCurrentElapsedTime();

‎GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ class GPUDebugTiming
6969
mRec->GPUFailedMsg(hipEventSynchronize((hipEvent_t)mDeviceTimers[1]));
7070
float v;
7171
mRec->GPUFailedMsg(hipEventElapsedTime(&v, (hipEvent_t)mDeviceTimers[0], (hipEvent_t)mDeviceTimers[1]));
72-
mXYZ.t = v * 1.e-3;
72+
mXYZ.t = v * 1.e-3f;
7373
} else {
7474
mRec->GPUFailedMsg(hipStreamSynchronize(mStreams[mXYZ.x.stream]));
7575
mXYZ.t = mTimer.GetCurrentElapsedTime();
@@ -127,7 +127,7 @@ void GPUReconstructionHIPBackend::runKernelBackendInternal(krnlSetup& _xyz, cons
127127
GPUFailedMsg(hipEventSynchronize((hipEvent_t)mDebugEvents->DebugStop));
128128
float v;
129129
GPUFailedMsg(hipEventElapsedTime(&v, (hipEvent_t)mDebugEvents->DebugStart, (hipEvent_t)mDebugEvents->DebugStop));
130-
_xyz.t = v * 1.e-3;
130+
_xyz.t = v * 1.e-3f;
131131
} else {
132132
backendInternal<T, I>::runKernelBackendMacro(_xyz, this, nullptr, nullptr, args...);
133133
}

‎GPU/GPUTracking/Base/hip/test/testGPUsortHIP.hip.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -52,7 +52,7 @@ struct TestEnvironment {
5252

5353
// create an array of unordered floats with negative and positive values
5454
for (size_t i = 0; i < size; i++) {
55-
data[i] = size / 2.0 - i;
55+
data[i] = size / 2.0f - i;
5656
}
5757
// create copy
5858
std::memcpy(sorted.data(), data, size * sizeof(float));

‎GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCLInternals.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -199,7 +199,7 @@ int GPUReconstructionOCL::runKernelBackendCommon(krnlSetup& _xyz, K& k, const Ar
199199
GPUFailedMsg(clWaitForEvents(1, evr));
200200
GPUFailedMsg(clGetEventProfilingInfo(*evr, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, nullptr));
201201
GPUFailedMsg(clGetEventProfilingInfo(*evr, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, nullptr));
202-
_xyz.t = (time_end - time_start) * 1.e-9;
202+
_xyz.t = (time_end - time_start) * 1.e-9f;
203203
if (tmpEvent) {
204204
GPUFailedMsg(clReleaseEvent(ev));
205205
}

‎GPU/GPUTracking/Base/opencl2/GPUReconstructionOCL2.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,7 @@ int GPUReconstructionOCL2Backend::GetOCLPrograms()
6464
const char* ocl_flags = GPUCA_M_STR(OCL_FLAGS);
6565

6666
#ifdef OPENCL2_ENABLED_SPIRV // clang-format off
67-
if (ver >= 2.2) {
67+
if (ver >= 2.2f) {
6868
GPUInfo("Reading OpenCL program from SPIR-V IL (Platform version %f)", ver);
6969
mInternals->program = clCreateProgramWithIL(mInternals->context, _binary_GPUReconstructionOCL2Code_spirv_start, _binary_GPUReconstructionOCL2Code_spirv_len, &ocl_error);
7070
ocl_flags = "";

‎GPU/GPUTracking/DataCompression/AliHLTTPCClusterStatComponent.cxx‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -172,7 +172,7 @@ void AliHLTTPCClusterStatComponent::TransformReverse(int slice, int row, float y
172172
padWidth = param->GetPadPitchWidth(sector);
173173
}
174174

175-
padtime[0] = y * sign / padWidth + 0.5 * maxPad;
175+
padtime[0] = y * sign / padWidth + 0.5f * maxPad;
176176

177177
float xyzGlobal[2] = {param->GetPadRowRadii(sector, sectorrow), y};
178178
AliHLTTPCGeometry::Local2Global(xyzGlobal, slice);
@@ -209,7 +209,7 @@ void AliHLTTPCClusterStatComponent::TransformForward(int slice, int row, float p
209209
}
210210

211211
xyz[0] = param->GetPadRowRadii(sector, sectorrow);
212-
xyz[1] = (pad - 0.5 * maxPad) * padWidth * sign;
212+
xyz[1] = (pad - 0.5f * maxPad) * padWidth * sign;
213213

214214
float xyzGlobal[2] = {xyz[0], xyz[1]};
215215
AliHLTTPCGeometry::Local2Global(xyzGlobal, slice);
@@ -364,8 +364,8 @@ int AliHLTTPCClusterStatComponent::DoEvent(const AliHLTComponentEventData& evtDa
364364

365365
int padrow = AliHLTTPCGeometry::GetFirstRow(patch) + cluster.GetPadRow();
366366
float x = AliHLTTPCGeometry::Row2X(padrow);
367-
float y = 0.0;
368-
float z = 0.0;
367+
float y = 0.0f;
368+
float z = 0.0f;
369369

370370
float xyz[3];
371371
if (1) // Use forward (exact reverse-reverse) transformation of raw cluster (track fit in distorted coordinates)

0 commit comments

Comments
 (0)