From bed8fb6f3007d82f2251285f971030e1d3e82a1c Mon Sep 17 00:00:00 2001 From: Christian Heinemann Date: Sat, 12 Oct 2024 22:12:41 +0200 Subject: [PATCH] macros corrected --- source/EngineGpuKernels/GenomeDecoder.cuh | 16 +++++++++++----- source/EngineGpuKernels/Macros.cuh | 8 ++++---- 2 files changed, 15 insertions(+), 9 deletions(-) diff --git a/source/EngineGpuKernels/GenomeDecoder.cuh b/source/EngineGpuKernels/GenomeDecoder.cuh index d06394db9..7048fd77f 100644 --- a/source/EngineGpuKernels/GenomeDecoder.cuh +++ b/source/EngineGpuKernels/GenomeDecoder.cuh @@ -63,6 +63,7 @@ public: __inline__ __device__ static void setNextConstructionAngle2(uint8_t* genome, int nodeAddress, uint8_t angle); __inline__ __device__ static void setNextConstructorSeparation(uint8_t* genome, int nodeAddress, bool separation); __inline__ __device__ static void setNextConstructorNumBranches(uint8_t* genome, int nodeAddress, int numBranches); + __inline__ __device__ static void setNextConstructorNumRepetitions(uint8_t* genome, int nodeAddress, int numRepetitions); __inline__ __device__ static bool containsSectionSelfReplication(uint8_t* genome, int genomeSize); __inline__ __device__ static int getNodeAddressForSelfReplication(uint8_t* genome, int genomeSize, bool& containsSelfReplicator); __inline__ __device__ static void setRandomCellFunctionData(SimulationData& data, uint8_t* genome, int nodeAddress, CellFunction const& cellFunction, bool makeSelfCopy, int subGenomeSize); @@ -116,7 +117,7 @@ __inline__ __device__ void GenomeDecoder::executeForEachNode(uint8_t* genome, in template __inline__ __device__ void GenomeDecoder::executeForEachNodeRecursively(uint8_t* genome, int genomeSize, bool includedSeparatedParts, bool countBranches, Func func) { - CHECK(genomeSize >= Const::GenomeHeaderSize) + CUDA_CHECK(genomeSize >= Const::GenomeHeaderSize) int subGenomeEndAddresses[MAX_SUBGENOME_RECURSION_DEPTH]; int subGenomeNumRepetitions[MAX_SUBGENOME_RECURSION_DEPTH + 1]; @@ -194,7 +195,7 @@ __inline__ __device__ int GenomeDecoder::getRandomGenomeNodeAddress( if (numSubGenomesSizeIndices) { *numSubGenomesSizeIndices = 0; } - CHECK(genomeSize >= Const::GenomeHeaderSize) + CUDA_CHECK(genomeSize >= Const::GenomeHeaderSize) if (genomeSize == Const::GenomeHeaderSize) { return Const::GenomeHeaderSize; @@ -329,8 +330,8 @@ __inline__ __device__ bool GenomeDecoder::hasInfiniteRepetitions(ConstructorFunc __inline__ __device__ bool GenomeDecoder::hasEmptyGenome(ConstructorFunction const& constructor) { if (constructor.genomeSize <= Const::GenomeHeaderSize) { - CHECK(constructor.genomeSize == Const::GenomeHeaderSize) - CHECK(constructor.genomeCurrentNodeIndex == 0) + CUDA_CHECK(constructor.genomeSize == Const::GenomeHeaderSize) + CUDA_CHECK(constructor.genomeCurrentNodeIndex == 0) return true; } return false; @@ -402,7 +403,7 @@ __inline__ __device__ bool GenomeDecoder::containsSelfReplication(ConstructorOrI __inline__ __device__ GenomeHeader GenomeDecoder::readGenomeHeader(ConstructorFunction const& constructor) { - CHECK(constructor.genomeSize >= Const::GenomeHeaderSize) + CUDA_CHECK(constructor.genomeSize >= Const::GenomeHeaderSize) GenomeHeader result; result.shape = constructor.genome[Const::GenomeHeaderShapePos] % ConstructionShape_Count; @@ -687,6 +688,11 @@ __inline__ __device__ void GenomeDecoder::setNextConstructorNumBranches(uint8_t* genome[nodeAddress + Const::CellBasicBytes + Const::ConstructorFixedBytes + 3 + Const::GenomeHeaderNumBranchesPos] = static_cast(numBranches); } +__inline__ __device__ void GenomeDecoder::setNextConstructorNumRepetitions(uint8_t* genome, int nodeAddress, int numRepetitions) +{ + genome[nodeAddress + Const::CellBasicBytes + Const::ConstructorFixedBytes + 3 + Const::GenomeHeaderNumRepetitionsPos] = static_cast(numRepetitions); +} + __inline__ __device__ int GenomeDecoder::getNextSubGenomeSize(uint8_t* genome, int genomeSize, int nodeAddress) { auto cellFunction = getNextCellFunctionType(genome, nodeAddress); diff --git a/source/EngineGpuKernels/Macros.cuh b/source/EngineGpuKernels/Macros.cuh index ed549c85d..629525f32 100644 --- a/source/EngineGpuKernels/Macros.cuh +++ b/source/EngineGpuKernels/Macros.cuh @@ -62,7 +62,7 @@ void checkAndThrowError(T result, char const *const func, const char *const file #define NEAR_ZERO 0.00001f -#define CHECK(condition) \ +#define CUDA_CHECK(condition) \ if (!(condition)) { \ printf("Check failed. File: %s, Line: %d\n", __FILE__, __LINE__); \ ABORT(); \ @@ -73,7 +73,7 @@ void checkAndThrowError(T result, char const *const func, const char *const file ABORT(); #define KERNEL_CALL(func, ...) \ - if (GlobalSettings::getInstance().isDebugMode()) { \ + if (GlobalSettings::get().isDebugMode()) { \ func<<>>(__VA_ARGS__); \ cudaDeviceSynchronize(); \ CHECK_FOR_CUDA_ERROR(cudaGetLastError()); \ @@ -83,7 +83,7 @@ void checkAndThrowError(T result, char const *const func, const char *const file } #define KERNEL_CALL_1_1(func, ...) \ - if (GlobalSettings::getInstance().isDebugMode()) { \ + if (GlobalSettings::get().isDebugMode()) { \ func<<<1, 1>>>(__VA_ARGS__); \ cudaDeviceSynchronize(); \ CHECK_FOR_CUDA_ERROR(cudaGetLastError()); \ @@ -92,7 +92,7 @@ void checkAndThrowError(T result, char const *const func, const char *const file } #define KERNEL_CALL_MOD(func, threadsPerBlock, ...) \ - if (GlobalSettings::getInstance().isDebugMode()) { \ + if (GlobalSettings::get().isDebugMode()) { \ func<<>>(__VA_ARGS__); \ cudaDeviceSynchronize(); \ CHECK_FOR_CUDA_ERROR(cudaGetLastError()); \