More CHECK_ERR cleanup

Author: davschneller

algorithms/Common.h

@@ -61,12 +61,9 @@ int blockcount(F&& func, int blocksize) {
   int smCount = 0;
   int blocksPerSM = 0;
   APIWRAP(cudaGetDevice(&device));
-  CHECK_ERR;
   APIWRAP(cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, device));
-  CHECK_ERR;
   APIWRAP(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
       &blocksPerSM, std::forward<F>(func), blocksize, 0));
-  CHECK_ERR;
   return smCount * blocksPerSM;
 }
 
@@ -90,12 +87,9 @@ int blockcount(F&& func, int blocksize) {
   int smCount = 0;
   int blocksPerSM = 0;
   APIWRAP(hipGetDevice(&device));
-  CHECK_ERR;
   APIWRAP(hipDeviceGetAttribute(&smCount, hipDeviceAttributeMultiprocessorCount, device));
-  CHECK_ERR;
   APIWRAP(hipOccupancyMaxActiveBlocksPerMultiprocessor(
       &blocksPerSM, std::forward<F>(func), blocksize, 0));
-  CHECK_ERR;
   return smCount * blocksPerSM;
 }
 

interfaces/cuda/Events.cu

@@ -20,7 +20,6 @@ void* ConcreteAPI::createEvent(bool withTiming) {
   } else {
     APIWRAP(cudaEventCreateWithFlags(&event, cudaEventDisableTiming));
   }
-  CHECK_ERR;
   return static_cast<void*>(event);
 }
 
@@ -35,31 +34,29 @@ double ConcreteAPI::timespanEvents(void* eventPtrStart, void* eventPtrEnd) {
 void ConcreteAPI::destroyEvent(void* eventPtr) {
   cudaEvent_t event = static_cast<cudaEvent_t>(eventPtr);
   APIWRAP(cudaEventDestroy(event));
-  CHECK_ERR;
 }
 
 void ConcreteAPI::syncEventWithHost(void* eventPtr) {
   cudaEvent_t event = static_cast<cudaEvent_t>(eventPtr);
-  APIWRAP(cudaEventSynchronize(event));
+
   CHECK_ERR;
+
+  APIWRAP(cudaEventSynchronize(event));
 }
 
 bool ConcreteAPI::isEventCompleted(void* eventPtr) {
   cudaEvent_t event = static_cast<cudaEvent_t>(eventPtr);
   auto result = APIWRAPX(cudaEventQuery(event), {cudaErrorNotReady});
-  CHECK_ERR;
   return result == cudaSuccess;
 }
 
 void ConcreteAPI::recordEventOnHost(void* eventPtr) {
   cudaEvent_t event = static_cast<cudaEvent_t>(eventPtr);
   APIWRAP(cudaEventRecord(event));
-  CHECK_ERR;
 }
 
 void ConcreteAPI::recordEventOnStream(void* eventPtr, void* streamPtr) {
   cudaEvent_t event = static_cast<cudaEvent_t>(eventPtr);
   cudaStream_t stream = static_cast<cudaStream_t>(streamPtr);
   APIWRAP(cudaEventRecord(event, stream));
-  CHECK_ERR;
 }

interfaces/cuda/Memory.cu

@@ -79,7 +79,6 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) {
     allocationProperties[devPtr] = reinterpret_cast<void*>(new CUmemAllocationProp(prop));
   } else {
     APIWRAP(cudaMalloc(&devPtr, size));
-    CHECK_ERR;
   }
   statistics.allocatedMemBytes += size;
   memToSizeMap[devPtr] = size;
@@ -90,7 +89,6 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint)
   isFlagSet<DeviceSelected>(status);
   void* devPtr = nullptr;
   APIWRAP(cudaMallocManaged(&devPtr, size, cudaMemAttachGlobal));
-  CHECK_ERR;
 
   cudaMemLocation location{};
   if (hint == Destination::Host) {
@@ -114,7 +112,6 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint)
                         location.id
 #endif
                         ));
-  CHECK_ERR;
 
   statistics.allocatedMemBytes += size;
   statistics.allocatedUnifiedMemBytes += size;
@@ -127,7 +124,6 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint)
   void* devPtr = nullptr;
   const auto flag = hint == Destination::Host ? cudaHostAllocDefault : cudaHostAllocMapped;
   APIWRAP(cudaHostAlloc(&devPtr, size, flag));
-  CHECK_ERR;
   statistics.allocatedMemBytes += size;
   memToSizeMap[devPtr] = size;
   return devPtr;

interfaces/cuda/Streams.cu

@@ -58,8 +58,6 @@ bool ConcreteAPI::isStreamWorkDone(void* streamPtr) {
   isFlagSet<InterfaceInitialized>(status);
   cudaStream_t stream = static_cast<cudaStream_t>(streamPtr);
 
-  CHECK_ERR;
-
   const auto streamStatus = APIWRAPX(cudaStreamQuery(stream), {cudaErrorNotReady});
 
   return streamStatus == cudaSuccess;

interfaces/hip/Events.cpp

@@ -39,6 +39,9 @@ void ConcreteAPI::destroyEvent(void* eventPtr) {
 
 void ConcreteAPI::syncEventWithHost(void* eventPtr) {
   hipEvent_t event = static_cast<hipEvent_t>(eventPtr);
+
+  CHECK_ERR;
+
   APIWRAP(hipEventSynchronize(event));
 }
 

interfaces/hip/Streams.cpp

@@ -47,6 +47,9 @@ void ConcreteAPI::destroyGenericStream(void* streamPtr) {
 void ConcreteAPI::syncStreamWithHost(void* streamPtr) {
   isFlagSet<InterfaceInitialized>(status);
   hipStream_t stream = static_cast<hipStream_t>(streamPtr);
+
+  CHECK_ERR;
+
   APIWRAP(hipStreamSynchronize(stream));
 }