From 999f45fc11bc4fb4596ef0f83b1398bddc9b2535 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Wed, 3 Jul 2019 20:42:19 -0400 Subject: [PATCH 1/4] [hip] Move _criticalData of ihipStream_t class to private section and use criticalData() to access it (#1177) --- src/hip_hcc.cpp | 4 ++-- src/hip_hcc_internal.h | 4 +--- src/hip_stream.cpp | 2 +- 3 files changed, 4 insertions(+), 6 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 4edf9575be..c0a3827376 100644 --- a/src/hip_hcc.cpp +++ b/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( diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index 852e8985ad..a07ff0e7e7 100644 --- a/src/hip_hcc_internal.h +++ b/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. diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index b748749538..d70f860f90 100644 --- a/src/hip_stream.cpp +++ b/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); From 2a1b0ba27dfba57b46613d129046d5d710d22380 Mon Sep 17 00:00:00 2001 From: Jannik Date: Thu, 4 Jul 2019 02:43:22 +0200 Subject: [PATCH 2/4] Fix call to cuda library (#1197) * Remove flags parameter from hipOccupancyMaxPotentialBlockSize This commit makes the hipOccupancyMaxPotentialBlockSize method consistent with hcc path and the CUDA API. --- include/hip/nvcc_detail/hip_runtime_api.h | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index c1846c1b1e..aca04067e3 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -1216,11 +1216,9 @@ inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t template 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); } From b6b749f714f7a2e8c421c46e36d6bc51c57a0ddd Mon Sep 17 00:00:00 2001 From: kpyzhov Date: Wed, 3 Jul 2019 20:44:06 -0400 Subject: [PATCH 3/4] Change the type of hipTextureObject_t to pointer to a struct. (#1199) This is necessary to allow Runtime to perform required texture buffer handling. --- include/hip/hcc_detail/texture_types.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/texture_types.h b/include/hip/hcc_detail/texture_types.h index bc334de248..ec3c78532c 100644 --- a/include/hip/hcc_detail/texture_types.h +++ b/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 From a9870f0465ffb9562db4cb2744dd4dd3f7e55255 Mon Sep 17 00:00:00 2001 From: searlmc1 Date: Wed, 3 Jul 2019 17:44:26 -0700 Subject: [PATCH 4/4] Fix typo ( T in Vec3_cmp = needs to be int ). (#1203) Typo introduced here: commit 67abac1365c26fdf499e288ac7f3eec88b52c816 Author: Alex Voicu Date: Mon Jun 24 20:02:09 2019 -0500 Put 3-wide vector types on a ketogenic diet. (#1180) --- include/hip/hcc_detail/hip_vector_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 0d424965ca..2ce216e409 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/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 {