Skip to content

Commit

Permalink
macros corrected
Browse files Browse the repository at this point in the history
  • Loading branch information
chrxh committed Oct 12, 2024
1 parent 95701a5 commit bed8fb6
Show file tree
Hide file tree
Showing 2 changed files with 15 additions and 9 deletions.
16 changes: 11 additions & 5 deletions source/EngineGpuKernels/GenomeDecoder.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -116,7 +117,7 @@ __inline__ __device__ void GenomeDecoder::executeForEachNode(uint8_t* genome, in
template <typename Func>
__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];
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -687,6 +688,11 @@ __inline__ __device__ void GenomeDecoder::setNextConstructorNumBranches(uint8_t*
genome[nodeAddress + Const::CellBasicBytes + Const::ConstructorFixedBytes + 3 + Const::GenomeHeaderNumBranchesPos] = static_cast<uint8_t>(numBranches);
}

__inline__ __device__ void GenomeDecoder::setNextConstructorNumRepetitions(uint8_t* genome, int nodeAddress, int numRepetitions)
{
genome[nodeAddress + Const::CellBasicBytes + Const::ConstructorFixedBytes + 3 + Const::GenomeHeaderNumRepetitionsPos] = static_cast<uint8_t>(numRepetitions);
}

__inline__ __device__ int GenomeDecoder::getNextSubGenomeSize(uint8_t* genome, int genomeSize, int nodeAddress)
{
auto cellFunction = getNextCellFunctionType(genome, nodeAddress);
Expand Down
8 changes: 4 additions & 4 deletions source/EngineGpuKernels/Macros.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -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(); \
Expand All @@ -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<<<gpuSettings.numBlocks, 8>>>(__VA_ARGS__); \
cudaDeviceSynchronize(); \
CHECK_FOR_CUDA_ERROR(cudaGetLastError()); \
Expand All @@ -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()); \
Expand All @@ -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<<<gpuSettings.numBlocks, threadsPerBlock>>>(__VA_ARGS__); \
cudaDeviceSynchronize(); \
CHECK_FOR_CUDA_ERROR(cudaGetLastError()); \
Expand Down

0 comments on commit bed8fb6

Please sign in to comment.