temporary fixes

Author: AtlantaPepsi

src/client/Presets/Help.hpp

@@ -113,11 +113,11 @@ int HelpPreset([[maybe_unused]] EnvVars&          ev,
   printf("## Single DMA executed Transfer between GPUs 0 and 1\n");
   printf("1 1 (G0->D0->G1)\n");
   printf("\n");
-  printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for tensor-op path\n");
+  printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for TDM path\n");
   printf("1 1 (G0->T0->G1)\n");
   printf("\n");
   printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for async load/store path\n");
-  printf("1 1 (G0->L0->G1)\n");
+  printf("1 1 (G0->A0->G1)\n");
   printf("\n");
   printf("## Copy 1Mb from GPU0 to GPU1 with 4 CUs, and 2Mb from GPU1 to GPU0 with 8 CUs\n");
   printf("-2 (G0->G0->G1 4 1M) (G1->G1->G0 8 2M)\n");

src/header/TransferBench.hpp

@@ -283,7 +283,7 @@ namespace TransferBench
 
   struct TdmOptions
   {
-    int blockSize    = 256; ///< Size of each threadblock (must be multiple of 64)
+    int blockSize    = 256; ///< Size of each threadblock (must be multiple of 32)
     int maxLDSBytes  = INT_MAX; ///< Maximum number of bytes of __shared__ memory to use
     int useHipEvents = 1; ///< Use HIP events for timing the TDM executor (0=CPU wall-clock)
 
@@ -638,6 +638,7 @@ namespace TransferBench
   // Enumerations
   #define hipDeviceAttributeClockRate                        cudaDevAttrClockRate
   #define hipDeviceAttributeMultiprocessorCount              cudaDevAttrMultiProcessorCount
+  #define hipDeviceAttributeMaxSharedMemoryPerBlock          cudaDevAttrMaxSharedMemoryPerBlock
   #define hipDeviceAttributeWarpSize                         cudaDevAttrWarpSize
   #define hipErrorPeerAccessAlreadyEnabled                   cudaErrorPeerAccessAlreadyEnabled
   #define hipFuncCachePreferShared                           cudaFuncCachePreferShared
@@ -667,6 +668,7 @@ namespace TransferBench
   #define hipGetDeviceCount                                  cudaGetDeviceCount
   #define hipGetDeviceProperties                             cudaGetDeviceProperties
   #define hipGetErrorString                                  cudaGetErrorString
+  #define hipGetLastError                                    cudaGetLastError
   #define hipHostFree                                        cudaFreeHost
   #define hipHostMalloc                                      cudaMallocHost
   #define hipMalloc                                          cudaMalloc
@@ -2396,12 +2398,18 @@ namespace {
           hasFatalError = true;
           break;
         }
+#if defined(__NVCC__)
+        errors.push_back({ERR_FATAL,
+                          "Transfer %d: GPU TDM kernel is not supported on NVIDIA hardware", i});
+        hasFatalError = true;
+#else
         if (!tdm::IsTdmCopySupported(t.exeDevice.exeIndex)) {
           errors.push_back({ERR_FATAL,
                             "Transfer %d: GPU TDM kernel requires gfx1250 hardware, but GPU %d is not supported",
                             i, t.exeDevice.exeIndex});
           hasFatalError = true;
         }
+#endif
         break;
       case EXE_GPU_GFX:
         if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) {
@@ -5833,6 +5841,7 @@ namespace {
     int subIterations = 0;
     while (1) {
       tdm::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes);
+      __syncthreads();
       if (++subIterations == numSubIterations) break;
     }
 
@@ -8407,6 +8416,7 @@ namespace {
 // Enumerations
 #undef hipDeviceAttributeClockRate
 #undef hipDeviceAttributeMultiprocessorCount
+#undef hipDeviceAttributeMaxSharedMemoryPerBlock
 #undef hipDeviceAttributeWarpSize
 #undef hipErrorPeerAccessAlreadyEnabled
 #undef hipFuncCachePreferShared
@@ -8437,6 +8447,7 @@ namespace {
 #undef hipGetDeviceCount
 #undef hipGetDeviceProperties
 #undef hipGetErrorString
+#undef hipGetLastError
 #undef hipHostFree
 #undef hipHostMalloc
 #undef hipMalloc

src/header/tdmCopy.h

@@ -23,7 +23,7 @@ THE SOFTWARE.
 ///
 ///        This is introduced via __device__ level memcpy-like API which can be
 ///        either blocking or asynchronous, and utilize all warps, or a team of
-///        congtigous warps, to allow for other warps to do other tasks
+///        contiguous warps, to allow for other warps to do other tasks
 ///
 /// \par Quick start
 /// \code