From b3c16ea7b51e0617e995f449e5312afd881fe2f1 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sun, 20 Nov 2016 12:18:08 -0600 Subject: [PATCH] Fixed hipDeviceGetCacheConfig on nvcc path 1. Changed test macro to emit line numbers 2. Added getcacheconfig api test for nvcc path 3. Fixed hipFuncCache_t data type TODO: With this commit, right now there are 2 func cache datatypes a. hipFuncCache_t for runtime API b. hipFuncCache for driver API Map these to a single data type Change-Id: Ia47c9f5d7c2633638051bf17b1103048a1ede973 --- .../include/hip/hcc_detail/hip_runtime_api.h | 4 +-- .../include/hip/nvcc_detail/hip_runtime_api.h | 5 ++++ hipamd/src/hip_device.cpp | 2 +- .../src/nvcc/Device/hipDeviceGetAttribute.cpp | 5 ++-- .../nvcc/Device/hipDeviceGetCacheConfig.cpp | 28 +++++++++++++++++++ hipamd/tests/src/test_common.h | 2 +- 6 files changed, 40 insertions(+), 6 deletions(-) create mode 100644 hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 09825e141e..65bec1ff1d 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -112,12 +112,12 @@ enum hipLimit_t /** * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored. */ -typedef enum hipFuncCache { +typedef enum hipFuncCache_t { hipFuncCachePreferNone, ///< no preference for shared memory or L1 (default) hipFuncCachePreferShared, ///< prefer larger shared memory and smaller L1 cache hipFuncCachePreferL1, ///< prefer larger L1 cache and smaller shared memory hipFuncCachePreferEqual, ///< prefer equal size L1 cache and shared memory -} hipFuncCache; +} hipFuncCache_t; /** diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 0d15dfcb01..c05d04ee29 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -76,6 +76,7 @@ typedef cudaStream_t hipStream_t; typedef cudaIpcEventHandle_t hipIpcEventHandle_t; typedef cudaIpcMemHandle_t hipIpcMemHandle_t; typedef cudaLimit hipLimit_t; +typedef cudaFuncCache hipFuncCache_t; typedef CUcontext hipCtx_t; typedef CUsharedconfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; @@ -323,6 +324,10 @@ inline static hipError_t hipDeviceSynchronize() { return hipCUDAErrorTohipError(cudaDeviceSynchronize()); } +inline static hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *pCacheConfig) { + return hipCUDAErrorTohipError(cudaDeviceGetCacheConfig(pCacheConfig)); +} + inline static const char* hipGetErrorString(hipError_t error){ return cudaGetErrorString(hipErrorToCudaError(error)); } diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 29ab0805b8..116f4ade20 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -71,7 +71,7 @@ hipError_t hipGetDeviceCount(int *count) return e; } -hipError_t hipDeviceSetCacheConfig(hipFuncCache cacheConfig) +hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { HIP_INIT_API(cacheConfig); diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp index f0e56bf274..a7e1814468 100644 --- a/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp @@ -23,12 +23,13 @@ THE SOFTWARE. int main() { int val; - hipDeviceAttr_t attr = hipDevAttrMaxThreadsPerBlock; + hipDeviceAttribute_t attr = hipDeviceAttributeMaxThreadsPerBlock; ///< Maximum number of threads per block. + HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0)); HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0)); HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, -1)); HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, -1)); - attr = 91; + attr = hipDeviceAttribute_t(91); HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0)); HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0)); diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp new file mode 100644 index 0000000000..17f7ed7fbd --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp @@ -0,0 +1,28 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include"test_common.h" + +int main() +{ + hipFuncCache_t pCacheConfig; + HIP_PRINT_STATUS(hipDeviceGetCacheConfig(&pCacheConfig)); + HIP_PRINT_STATUS(hipDeviceGetCacheConfig(NULL)); +} diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index ee207db39d..90c130c56a 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -65,7 +65,7 @@ THE SOFTWARE. printf ("warn: TEST WARNING\n%s", KNRM );\ #define HIP_PRINT_STATUS(status) \ - std::cout<