Revert "Fix occupany APIs (#1560)"

This reverts commit af351d7e1b.

src/hip_module.cpp

@@ -1248,9 +1248,9 @@ void getGprsLdsUsage(hipFunction_t f, size_t* usedVGPRS, size_t* usedSGPRS, size
     }
 }
 
-hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, int* gridSize, int* blockSize,
-                                              hipFunction_t f, size_t dynamicSMemSize,
-                                              int blockSizeLimit)
+hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, uint32_t* gridSize, uint32_t* blockSize,
+                                              hipFunction_t f, size_t dynSharedMemPerBlk,
+                                              uint32_t blockSizeLimit)
 {
     using namespace hip_impl;
 
@@ -1331,7 +1331,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, int* gridSize, int*
         }
         else {
             size_t availableSharedMemPerCU = prop.maxSharedMemoryPerMultiProcessor;
-            size_t workgroupPerCU = availableSharedMemPerCU / (usedLDS + dynamicSMemSize);
+            size_t workgroupPerCU = availableSharedMemPerCU / (usedLDS + dynSharedMemPerBlk);
             wavefrontsLDS = min(workgroupPerCU, maxWorkgroupPerCU) * wavefrontsPerWG;
         }
 
@@ -1360,19 +1360,18 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, int* gridSize, int*
     return hipSuccess;
 }
 
-hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
-                                             const void* f, size_t dynamicSMemSize,
-                                             int blockSizeLimit)
+hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
+                                             hipFunction_t f, size_t dynSharedMemPerBlk,
+                                             uint32_t blockSizeLimit)
 {
-    HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, gridSize, blockSize, f, dynamicSMemSize, blockSizeLimit);
-    auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f),
-                                                   hip_impl::target_agent(0));
+    HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit);
+
     return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize(tls,
-        gridSize, blockSize, F, dynamicSMemSize, blockSizeLimit));
+        gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit));
 }
 
 hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
-   TlsData *tls, int* numBlocks, hipFunction_t f, int blockSize, size_t dynamicSMemSize)
+   TlsData *tls, uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk)
 {
     using namespace hip_impl;
 
@@ -1412,39 +1411,35 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
         : std::min(maxWavesPerSimd, availableSGPRs / usedSGPRS));
 
     // Calculate blocks occupancy per CU based on SGPR usage
-    *numBlocks = std::min(*numBlocks, (int) (sgprs_alu_occupancy / numWavefronts));
+    *numBlocks = std::min(*numBlocks, (uint32_t) (sgprs_alu_occupancy / numWavefronts));
 
-    size_t total_used_lds = usedLDS + dynamicSMemSize;
+    size_t total_used_lds = usedLDS + dynSharedMemPerBlk;
     if (total_used_lds != 0) {
       // Calculate LDS occupacy per CU. lds_per_cu / (static_lsd + dynamic_lds)
       size_t lds_occupancy = prop.maxSharedMemoryPerMultiProcessor / total_used_lds;
-      *numBlocks = std::min(*numBlocks, (int) lds_occupancy);
+      *numBlocks = std::min(*numBlocks, (uint32_t) lds_occupancy);
     }
 
     return hipSuccess;
 }
 
 hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
-   int* numBlocks, const void* f, int blockSize, size_t dynamicSMemSize)
+   uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk)
 {
-    HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynamicSMemSize);
+    HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk);
 
-    auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f),
-                                                   hip_impl::target_agent(0));
     return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
-        tls, numBlocks, F, blockSize, dynamicSMemSize));
+        tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
 }
 
 hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
-   int* numBlocks, const void* f, int  blockSize, size_t dynamicSMemSize,
+   uint32_t* numBlocks, hipFunction_t f, uint32_t  blockSize, size_t dynSharedMemPerBlk,
    unsigned int flags)
 {
-    HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynamicSMemSize, flags);
+    HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags);
 
-    auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f),
-                                                   hip_impl::target_agent(0));
     return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
-        tls, numBlocks, F, blockSize, dynamicSMemSize));
+        tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
 }
 
 hipError_t hipLaunchKernel(