Merge 'master' into 'amd-master'

Change-Id: Ie33c9467b47fc5d42da71b9bfc2932c26493f9f2

include/hip/hcc_detail/hip_vector_types.h

@@ -210,7 +210,7 @@ THE SOFTWARE.
                 return *this;
             }
 
-            using Vec3_cmp = T __NATIVE_VECTOR__(3, int);
+            using Vec3_cmp = int __NATIVE_VECTOR__(3, int);
             __host__ __device__
             Vec3_cmp operator==(const Native_vec_& x) const noexcept
             {

include/hip/hcc_detail/texture_types.h

@@ -45,7 +45,8 @@ THE SOFTWARE.
 /**
  * An opaque value that represents a hip texture object
  */
-typedef unsigned long long hipTextureObject_t;
+struct __hip_texture;
+typedef struct __hip_texture* hipTextureObject_t;
 
 /**
  * hip texture address modes

include/hip/nvcc_detail/hip_runtime_api.h

@@ -1216,11 +1216,9 @@ inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t
 template <class T>
 inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func,
                                                            size_t dynamicSMemSize = 0,
-                                                           int blockSizeLimit = 0,
-                                                           unsigned int flags = 0) {
+                                                           int blockSizeLimit = 0) {
     cudaError_t cerror;
-    cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize,
-                                                blockSizeLimit, flags);
+    cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit);
     return hipCUDAErrorTohipError(cerror);
 }
 

src/hip_hcc.cpp

@@ -1038,7 +1038,7 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
             }
         } else {
             if (waitThisStream) {
-                LockedAccessor_StreamCrit_t streamCrit(stream->_criticalData);
+                LockedAccessor_StreamCrit_t streamCrit(stream->criticalData());
 
                 // The last marker will provide appropriate visibility:
                 if (!streamCrit->_av.get_is_empty()) {
@@ -1056,7 +1056,7 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
 
     // Enqueue a barrier to wait on all the barriers we sent above:
     if (!HIP_SYNC_NULL_STREAM && !depOps.empty()) {
-        LockedAccessor_StreamCrit_t defaultStreamCrit(_defaultStream->_criticalData);
+        LockedAccessor_StreamCrit_t defaultStreamCrit(_defaultStream->criticalData());
         tprintf(DB_SYNC, "  null-stream wait on %zu non-empty streams. sync_host=%d\n",
                 depOps.size(), syncHost);
         hc::completion_future defaultCf = defaultStreamCrit->_av.create_blocking_marker(

src/hip_hcc_internal.h

@@ -609,12 +609,10 @@ class ihipStream_t {
 
     void addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes);
 
-
-   public:  // TODO - move private
+   private:  // Data
     // Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t
     ihipStreamCritical_t _criticalData;
 
-   private:  // Data
     std::mutex _hasQueueLock;
 
     ihipCtx_t* _ctx;  // parent context that owns this stream.

src/hip_stream.cpp

@@ -246,7 +246,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
 #if defined(__HCC__) && (__hcc_minor__ < 3)
         *priority = 0;
 #else
-        LockedAccessor_StreamCrit_t crit(stream->_criticalData);
+        LockedAccessor_StreamCrit_t crit(stream->criticalData());
         *priority = crit->_av.get_queue_priority();
 #endif
         return ihipLogStatus(hipSuccess);