Skip to content

Commit e398192

Browse files
davidrohrdavidrohr
committed
GPU: Use dummy member shadowing constexpr-replaced members to maintain same memory layout
1 parent 2ae22ff commit e398192

5 files changed

Lines changed: 15 additions & 53 deletions

File tree

‎GPU/GPUTracking/Base/GPUParam.cxx‎

Lines changed: 2 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -235,9 +235,7 @@ void GPUParam::LoadClusterErrors(bool Print)
235235

236236
void GPUParamRTC::setFrom(const GPUParam& param)
237237
{
238-
memcpy((char*)this + sizeof(gpu_rtc::GPUSettingsRec) + sizeof(gpu_rtc::GPUSettingsParam), (char*)&param + sizeof(GPUSettingsRec) + sizeof(GPUSettingsParam), sizeof(param) - sizeof(GPUSettingsRec) - sizeof(GPUSettingsParam));
239-
qConfigConvertRtc(this->rec, param.rec);
240-
qConfigConvertRtc(this->par, param.par);
238+
memcpy((char*)this, (char*)&param, sizeof(param));
241239
}
242240

243241
std::string GPUParamRTC::generateRTCCode(const GPUParam& param, bool useConstexpr)
@@ -249,10 +247,7 @@ std::string GPUParamRTC::generateRTCCode(const GPUParam& param, bool useConstexp
249247
qConfigPrintRtc(std::make_tuple(&param.rec, &param.par), useConstexpr);
250248
}
251249

252-
static_assert(alignof(GPUCA_NAMESPACE::gpu::GPUParam) == alignof(GPUCA_NAMESPACE::gpu::GPUSettingsRec));
253-
static_assert(alignof(GPUCA_NAMESPACE::gpu::GPUParam) == alignof(GPUCA_NAMESPACE::gpu::GPUSettingsParam));
254-
static_assert(sizeof(GPUCA_NAMESPACE::gpu::GPUParam) - sizeof(GPUCA_NAMESPACE::gpu::GPUParamRTC) == sizeof(GPUCA_NAMESPACE::gpu::GPUSettingsRec) + sizeof(GPUCA_NAMESPACE::gpu::GPUSettingsParam) - sizeof(GPUCA_NAMESPACE::gpu::gpu_rtc::GPUSettingsRec) - sizeof(GPUCA_NAMESPACE::gpu::gpu_rtc::GPUSettingsParam));
255-
static_assert(sizeof(GPUParam) % alignof(GPUConstantMem) == 0 && sizeof(GPUParamRTC) % alignof(GPUConstantMem) == 0, "Size of both GPUParam and of GPUParamRTC must be a multiple of the alignmeent of GPUConstantMem");
250+
static_assert(sizeof(GPUCA_NAMESPACE::gpu::GPUParam) == sizeof(GPUCA_NAMESPACE::gpu::GPUParamRTC), "RTC param size mismatch");
256251

257252
o2::base::Propagator* GPUParam::GetDefaultO2Propagator(bool useGPUField) const
258253
{

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu‎

Lines changed: 4 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -471,12 +471,9 @@ int GPUReconstructionCUDABackend::InitDevice_Runtime()
471471
}
472472
#else
473473
GPUFailedMsg(cudaMalloc(&devPtrConstantMem, gGPUConstantMemBufferSize));
474-
if (mProcessingSettings.enableRTC) {
475-
GPUFailedMsg(cudaMalloc(&devPtrConstantMemRTC, gGPUConstantMemBufferSize));
476-
}
474+
devPtrConstantMemRTC = devPtrConstantMem;
477475
#endif
478476
mDeviceConstantMem = (GPUConstantMem*)devPtrConstantMem;
479-
mDeviceConstantMem = (GPUConstantMem*)devPtrConstantMemRTC; // TODO: This is a hack, since the memory is sometimes addressed via the pointer from the kernel, should be made consistent!
480477
if (mProcessingSettings.enableRTC) {
481478
mDeviceConstantMemRTC.resize(1);
482479
mDeviceConstantMemRTC[0] = devPtrConstantMemRTC;
@@ -533,10 +530,6 @@ int GPUReconstructionCUDABackend::ExitDevice_Runtime()
533530
GPUFailedMsgI(cudaFree(mDeviceMemoryBase));
534531
#ifdef GPUCA_NO_CONSTANT_MEMORY
535532
GPUFailedMsgI(cudaFree(mDeviceConstantMem));
536-
for (unsigned int i = 0; i < mDeviceConstantMemRTC.size(); i++) {
537-
GPUFailedMsgI(cudaFree(mDeviceConstantMemRTC[i]));
538-
}
539-
mDeviceConstantMemRTC.clear();
540533
#endif
541534

542535
for (int i = 0; i < mNStreams; i++) {
@@ -611,20 +604,10 @@ size_t GPUReconstructionCUDABackend::WriteToConstantMemory(size_t offset, const
611604
int iLast = 1 + mDeviceConstantMemRTC.size();
612605
std::unique_ptr<GPUParamRTC> tmpParam;
613606
for (int i = iFirst; i < iLast; i++) {
614-
if (i == 1) { // First time we copy into the RTC constant mem, need to prepare
615-
if (offset < sizeof(GPUParam) && (offset != 0 || size > sizeof(GPUParam))) {
616-
throw std::runtime_error("Invalid write to constant memory, crossing GPUParam border");
617-
}
618-
if (offset == 0) {
619-
tmpParam.reset(new GPUParamRTC);
620-
tmpParam->setFrom(*(GPUParam*)src);
621-
src = tmpParam.get();
622-
size = sizeof(*tmpParam);
623-
} else {
624-
offset = offset - sizeof(GPUParam) + sizeof(GPUParamRTC);
625-
}
626-
}
627607
void* basePtr = i ? mDeviceConstantMemRTC[i - 1] : mDeviceConstantMem;
608+
if (i && basePtr == (void*)mDeviceConstantMem) {
609+
continue;
610+
}
628611
if (stream == -1) {
629612
GPUFailedMsg(cudaMemcpy(((char*)basePtr) + offset, src, size, cudaMemcpyHostToDevice));
630613
} else {

‎GPU/GPUTracking/Definitions/GPUSettingsList.h‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -88,7 +88,6 @@ AddOptionRTC(trdStopTrkAfterNMissLy, unsigned char, 6, "", 0, "Abandon track fol
8888
AddOptionRTC(trackingRefitGPUModel, char, 1, "", 0, "Use GPU track model for the Global Track Refit")
8989
AddOptionRTC(dropSecondaryLegsInOutput, char, 1, "", 0, "Do not store secondary legs of looping track in TrackTPC")
9090
AddCustomCPP(void SetMinTrackPt(float v) { MaxTrackQPt = v > 0.001 ? (1. / v) : (1. / 0.001); })
91-
AddVariable(dummyRTC, void*, nullptr) // Ensure non empty struct and proper alignment even if all normal members are constexpr
9291
AddHelp("help", 'h')
9392
EndConfig()
9493

@@ -368,7 +367,6 @@ AddVariableRTC(dodEdx, char, 0) // Do dEdx computation
368367
AddVariableRTC(earlyTpcTransform, char, 0) // do Early TPC transformation
369368
AddVariableRTC(debugLevel, char, 0) // Debug level
370369
AddVariableRTC(continuousMaxTimeBin, int, 0) // Max time bin for continuous tracking
371-
AddVariable(dummyRTC, void*, nullptr) // Ensure non empty struct and proper alignment even if all normal members are constexpr
372370
EndConfig()
373371

374372
EndNamespace() // gpu

‎GPU/GPUTracking/utils/qconfig.h‎

Lines changed: 5 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -248,6 +248,7 @@ enum qConfigRetVal { qcrOK = 0,
248248
#define AddVariableRTC(name, type, default) \
249249
if (useConstexpr) { \
250250
out << "static constexpr " << qon_mxstr(type) << " " << qon_mxstr(name) << " = " << qConfig::print_type(std::get<const qConfigCurrentType*>(tSrc)->name) << ";\n"; \
251+
out << qon_mxstr(type) << " " << qon_mxstr(qon_mxcat(_dummy_, name)) << ";\n"; \
251252
} else { \
252253
AddOption(name, type, default, optname, optnameshort, help); \
253254
}
@@ -266,27 +267,6 @@ enum qConfigRetVal { qcrOK = 0,
266267
#define AddOptionSet(...)
267268

268269
// End QCONFIG_PRINT_RTC
269-
#elif defined(QCONFIG_CONVERT_RTC)
270-
#define AddOption(name, type, default, optname, optnameshort, help, ...) out.name = in.name;
271-
#define AddVariable(name, type, default) out.name = in.name;
272-
#define AddOptionArray(name, type, count, default, optname, optnameshort, help, ...) \
273-
for (unsigned int i = 0; i < count; i++) \
274-
out.name[i] = in.name[i];
275-
#define AddOptionVec(name, type, optname, optnameshort, help, ...) \
276-
for (unsigned int i = 0; i < in.name.size(); i++) \
277-
out.name[i] = in.name[i];
278-
#define AddOptionRTC(name, type, default, optname, optnameshort, help, ...)
279-
#define AddVariableRTC(name, type, default)
280-
#define BeginConfig(name, instance) \
281-
template <class T> \
282-
void qConfigConvertRtc(T& out, const name& in) \
283-
{
284-
#define BeginSubConfig(name, instance, parent, preoptname, preoptnameshort, descr) BeginConfig(name, instance)
285-
#define EndConfig() }
286-
#define AddSubConfig(name, instance)
287-
#define AddOptionSet(...)
288-
289-
// End QCONFIG_CONVERT_RTC
290270
#else // Define structures
291271
#if defined(QCONFIG_HEADER_GUARD) && !defined(QCONFIG_GENRTC)
292272
#define QCONFIG_HEADER_GUARD_NO_INCLUDE
@@ -311,8 +291,10 @@ enum qConfigRetVal { qcrOK = 0,
311291
#define AddOptionVec(name, type, optname, optnameshort, help, ...) void* name[sizeof(std::vector<type>) / sizeof(void*)];
312292
#endif
313293
#ifdef QCONFIG_GENRTC
314-
#define AddOptionRTC(name, type, default, optname, optnameshort, help, ...) static constexpr type name = default;
315-
#define AddVariableRTC(name, type, default) static constexpr type name = default;
294+
#define AddVariableRTC(name, type, default) \
295+
static constexpr type name = default; \
296+
type _dummy_##name = default;
297+
#define AddOptionRTC(name, type, default, optname, optnameshort, help, ...) AddVariableRTC(name, type, default)
316298
#else
317299
#define AddCustomCPP(...) __VA_ARGS__
318300
#endif

‎GPU/GPUTracking/utils/qconfigrtc.h‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,10 @@
2121
#define qon_mstr(a) #a
2222
#define qon_mxstr(a) qon_mstr(a)
2323
#endif
24+
#ifndef qon_mxcat
25+
#define qon_mcat(a, b) a##b
26+
#define qon_mxcat(a, b) qon_mcat(a, b)
27+
#endif
2428

2529
template <class T>
2630
static std::string qConfigPrintRtc(const T& tSrc, bool useConstexpr)

0 commit comments

Comments
 (0)