diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 07cc22f4f9..64fb0abcd1 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1151,8 +1151,6 @@ INLINE ihipDevice_t *ihipGetDevice(int deviceId) } - - //--- // Get the stream to use for a command submission. // @@ -1162,8 +1160,10 @@ inline hipStream_t ihipSyncAndResolveStream(hipStream_t stream) { if (stream == hipStreamNull ) { ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->syncDefaultStream(false); +#ifndef HIP_API_PER_THREAD_DEFAULT_STREAM + device->syncDefaultStream(false); +#endif return device->_default_stream; } else { // Have to wait for legacy default stream to be empty: diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index d6031e679e..84b1dfee0b 100644 --- a/tests/src/CMakeLists.txt +++ b/tests/src/CMakeLists.txt @@ -109,6 +109,8 @@ macro (make_test_matches exe match_string) ) endmacro() +make_hip_executable (hipAPIStreamEnable hipAPIStreamEnable.cpp) +make_hip_executable (hipAPIStreamDisable hipAPIStreamDisable.cpp) make_hip_executable (hip_ballot hip_ballot.cpp) make_hip_executable (hip_anyall hip_anyall.cpp) make_hip_executable (hip_popc hip_popc.cpp) @@ -138,7 +140,6 @@ make_hip_executable (hipStreamL5 hipStreamL5.cpp) make_hip_executable (hipHostGetFlags hipHostGetFlags.cpp) make_hip_executable (hipHostRegister hipHostRegister.cpp) make_hip_executable (hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp) - make_test(hip_ballot " " ) make_test(hip_anyall " " ) make_test(hip_popc " " ) @@ -159,6 +160,7 @@ make_test(hipMemcpy_simple " " ) make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 ) make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 ) make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 ) + make_test(hipHostAlloc " ") make_test(hipMemcpyAsync " " ) make_test(hipHostGetFlags " ") @@ -166,4 +168,6 @@ make_test(hipHcc " " ) make_test(hipHostRegister " ") make_test(hipStreamL5 " ") make_test(hipRandomMemcpyAsync " ") +make_test(hipAPIStreamEnable " ") +make_test(hipAPIStreamDisable " ") make_hipify_test(specialFunc.cu ) diff --git a/tests/src/hipAPIStreamDisable.cpp b/tests/src/hipAPIStreamDisable.cpp new file mode 100644 index 0000000000..a419478458 --- /dev/null +++ b/tests/src/hipAPIStreamDisable.cpp @@ -0,0 +1,61 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include"test_common.h" + +const int NN = 1 << 21; + +__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){ + int tid = hipThreadIdx_x; + if(tid < 1){ + for(int i=0;i +#include"test_common.h" + +const int NN = 1 << 21; + +__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){ + int tid = hipThreadIdx_x; + if(tid < 1){ + for(int i=0;i