From 4353cefd6ecdbf45a801c3b586703b20a50294dc Mon Sep 17 00:00:00 2001 From: pensun Date: Tue, 16 Aug 2016 15:28:42 -0500 Subject: [PATCH] add occupancy support for NV path; fix hipPeekAtLastError on HCC path Change-Id: I26b0e1875c19d7c636ffcc18f1738926572ded81 [ROCm/clr commit: e70cfd3582a780d931571d52358ec8b6eca3199c] --- .../include/nvcc_detail/hip_runtime_api.h | 29 +++++++++++++++++-- projects/clr/hipamd/src/hip_error.cpp | 5 ++-- 2 files changed, 29 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h index 6ea49fb3a1..a51bca2d02 100644 --- a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h @@ -82,7 +82,7 @@ switch(cuError) { case cudaErrorHostMemoryAlreadyRegistered : return hipErrorHostMemoryAlreadyRegistered ; case cudaErrorHostMemoryNotRegistered : return hipErrorHostMemoryNotRegistered ; default : return hipErrorUnknown; // Note - translated error. -}; +}; } // TODO match the error enum names of hip and cuda @@ -108,7 +108,7 @@ switch(hError) { case hipErrorHostMemoryNotRegistered : return cudaErrorHostMemoryNotRegistered ; case hipErrorTbd : return cudaErrorUnknown; // Note - translated error. default : return cudaErrorUnknown; // Note - translated error. -} +} } inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) { @@ -347,6 +347,31 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att return hipCUDAErrorTohipError(cerror); } +template +inline static hipError_t hipOccupancyMaxPotentialBlockSize( + int *minGridSize, + int *blockSize, + T func, + size_t dynamicSMemSize = 0, + int blockSizeLimit = 0, + unsigned int flags = 0 + ){ + cudaError_t cerror; + cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit, flags); + return hipCUDAErrorTohipError(cerror); +} + +inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( + int *numBlocks, + const void* func, + int blockSize, + size_t dynamicSMemSize + ) +{ + cudaError_t cerror; + cerror = cudaOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, func, blockSize, dynamicSMemSize); + return hipCUDAErrorTohipError(cerror); +} inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr){ cudaPointerAttributes cPA; diff --git a/projects/clr/hipamd/src/hip_error.cpp b/projects/clr/hipamd/src/hip_error.cpp index 7c723b1aa2..d9c6dd9aa9 100644 --- a/projects/clr/hipamd/src/hip_error.cpp +++ b/projects/clr/hipamd/src/hip_error.cpp @@ -40,12 +40,11 @@ hipError_t hipGetLastError() //--- -hipError_t hipPeakAtLastError() +hipError_t hipPeekAtLastError() { HIP_INIT_API(); - - // peak at last error, but don't reset it. + // peek at last error, but don't reset it. return ihipLogStatus(tls_lastHipError); }