Merge 'amd-master-next' into 'amd-npi-next'

Change-Id: I56b40b684a8be0dfbf330d0d1ed1ad703fde8417

docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md

@@ -103,7 +103,7 @@
 |   **CUDA**                                                |   **HIP**                             |**CUDA version\***|
 |-----------------------------------------------------------|---------------------------------------|:----------------:|
 | `cudaFuncGetAttributes`                                   |`hipFuncGetAttributes`                 |
-| `cudaFuncSetAttribute`                                    |                                       | 9.0              |
+| `cudaFuncSetAttribute`                                    |`hipFuncSetAttribute`                  | 9.0              |
 | `cudaFuncSetCacheConfig`                                  |`hipFuncSetCacheConfig`                |
 | `cudaFuncSetSharedMemConfig`                              |                                       |
 | `cudaGetParameterBuffer`                                  |                                       |

include/hip/hcc_detail/hip_runtime_api.h

@@ -290,6 +290,14 @@ typedef enum hipJitOption {
     hipJitOptionNumOptions
 } hipJitOption;
 
+/**
+ * @warning On AMD devices and some Nvidia devices, these hints and controls are ignored.
+ */
+typedef enum hipFuncAttribute {
+    hipFuncAttributeMaxDynamicSharedMemorySize = 8,
+    hipFuncAttributePreferredSharedMemoryCarveout = 9,
+    hipFuncAttributeMax
+} hipFuncAttribute;
 
 /**
  * @warning On AMD devices and some Nvidia devices, these hints and controls are ignored.
@@ -528,6 +536,21 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig);
 hipError_t hipDeviceGetLimit(size_t* pValue, enum hipLimit_t limit);
 
 
+/**
+ * @brief Set attribute for a specific function
+ *
+ * @param [in] func;
+ * @param [in] attr;
+ * @param [in] value;
+ *
+ * @returns #hipSuccess, #hipErrorInvalidDeviceFunction, #hipErrorInvalidValue
+ *
+ * Note: AMD devices and some Nvidia GPUS do not support shared cache banking, and the hint is
+ * ignored on those architectures.
+ *
+ */
+hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value);
+
 /**
  * @brief Set Cache configuration for a specific function
  *

include/hip/nvcc_detail/hip_runtime_api.h

@@ -185,6 +185,7 @@ typedef cudaStream_t hipStream_t;
 typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
 typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
 typedef enum cudaLimit hipLimit_t;
+typedef enum cudaFuncAttribute hipFuncAttribute;
 typedef enum cudaFuncCache hipFuncCache_t;
 typedef CUcontext hipCtx_t;
 typedef enum cudaSharedMemConfig hipSharedMemConfig;
@@ -1007,6 +1008,10 @@ inline static hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* pCacheConfig) {
     return hipCUDAErrorTohipError(cudaDeviceGetCacheConfig(pCacheConfig));
 }
 
+inline static hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value) {
+    return hipCUDAErrorTohipError(cudaFuncSetAttribute(func, attr, value));
+}
+
 inline static hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) {
     return hipCUDAErrorTohipError(cudaDeviceSetCacheConfig(cacheConfig));
 }

rocclr/hip_event.cpp

@@ -143,7 +143,13 @@ void Event::addMarker(amd::HostQueue* queue, amd::Command* command, bool record)
   if (queue->properties().test(CL_QUEUE_PROFILING_ENABLE)) {
     if (command == nullptr) {
       command = queue->getLastQueuedCommand(true);
-      if (command == nullptr) {
+      if ((command == nullptr) || (command->type() == 0)) {
+        // if lastEnqueuedCommand is user invisible command(command->type() == 0),
+        // which is only used for sync, create a new amd:Marker
+        // and release() lastEnqueuedCommand
+        if (command != nullptr) {
+          command->release();
+        }
         command = new amd::Marker(*queue, kMarkerDisableFlush);
         command->enqueue();
       }

rocclr/hip_hcc.def.in

@@ -55,6 +55,7 @@ hipExtModuleLaunchKernel
 hipExtLaunchKernel
 hipFree
 hipFreeArray
+hipFuncSetAttribute
 hipFuncSetCacheConfig
 hipGetDevice
 hipGetDeviceCount

rocclr/hip_hcc.map.in

@@ -56,6 +56,7 @@ global:
     hipExtLaunchKernel;
     hipFree;
     hipFreeArray;
+    hipFuncSetAttribute;
     hipFuncSetCacheConfig;
     hipGetDevice;
     hipGetDeviceCount;

rocclr/hip_module.cpp

@@ -193,6 +193,14 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
   HIP_RETURN(hipSuccess);
 }
 
+hipError_t hipFuncSetAttribute ( const void* func, hipFuncAttribute attr, int value ) {
+  HIP_INIT_API(hipFuncSetAttribute, func, attr, value);
+
+  // No way to set function attribute yet.
+
+  HIP_RETURN(hipSuccess);
+}
+
 hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                  uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                  uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,

samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp

@@ -29,6 +29,7 @@ THE SOFTWARE.
 #include <thread>
 #include <future>
 #include <functional>
+#include <vector>
 
 #define NUM_GROUPS 1
 #define GROUP_SIZE 1

src/hip_module.cpp

@@ -1408,6 +1408,14 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
     return ihipLogStatus(hipSuccess);
 }
 
+hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value) {
+    HIP_INIT_API(hipFuncSetCacheConfig, func, attr, value);
+
+    // Nop, AMD does not support setting shared memory size for function.
+
+    return ihipLogStatus(hipSuccess);
+}
+
 hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc)
 {
     HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc);