From e65d4860f3d012a5f16af5f2a24149dcf69c172e Mon Sep 17 00:00:00 2001 From: Christian Heinemann Date: Tue, 21 May 2024 23:19:36 +0200 Subject: [PATCH] Revert "better solution for GCC compile error" This reverts commit a1546a5c724695fa1077263f7515fa5b6c398028. --- source/EngineGpuKernels/Base.cuh | 8 +++----- source/EngineGpuKernels/DensityMap.cuh | 6 +++--- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/source/EngineGpuKernels/Base.cuh b/source/EngineGpuKernels/Base.cuh index 19ddda940..b43b4e08e 100644 --- a/source/EngineGpuKernels/Base.cuh +++ b/source/EngineGpuKernels/Base.cuh @@ -122,13 +122,11 @@ __device__ __inline__ T alienAtomicRead(T* const& address) // // Due to this, we need an ugly type casting workaround here // -template -__device__ __inline__ U alienAtomicAdd64(S* address, T const& value) +template +__device__ __inline__ T alienAtomicAdd64(T* address, T const& value) { static_assert(sizeof(unsigned long long) == sizeof(T)); - static_assert(sizeof(unsigned long long) == sizeof(S)); - static_assert(sizeof(unsigned long long) == sizeof(U)); - return reinterpret_cast(atomicAdd(reinterpret_cast(address), static_cast(value))); + return atomicAdd(reinterpret_cast(address), static_cast(value)); } template diff --git a/source/EngineGpuKernels/DensityMap.cuh b/source/EngineGpuKernels/DensityMap.cuh index c99c5f3f7..b56b0f73d 100644 --- a/source/EngineGpuKernels/DensityMap.cuh +++ b/source/EngineGpuKernels/DensityMap.cuh @@ -64,13 +64,13 @@ public: auto index = toInt(cell->pos.x) / _slotSize + toInt(cell->pos.y) / _slotSize * _densityMapSize.x; if (index >= 0 && index < _densityMapSize.x * _densityMapSize.y) { auto color = calcMod(cell->color, MAX_COLORS); - alienAtomicAdd64(&_colorDensityMap[index], (1ull << (color * 8)) | (1ull << 56)); + alienAtomicAdd64(&_colorDensityMap[index], static_cast((1ull << (color * 8)) | (1ull << 56))); if (cell->mutationId != 0) { uint64_t bucket = calcBucket(cell->mutationId, timestep); - alienAtomicAdd64(&_otherMutantDensityMap[index], 0x0101010101010101ull ^ (1ull << (bucket * 8))); + alienAtomicAdd64(&_otherMutantDensityMap[index], static_cast(0x0101010101010101ull ^ (1ull << (bucket * 8)))); } else { - alienAtomicAdd64(&_otherMutantDensityMap[index], 1ull); + alienAtomicAdd64(&_otherMutantDensityMap[index], static_cast(1ull)); } } }