fixed directed tests fail when hcc bumped to 3.0 (#1678)

Handled the HCC version check appropriately as few of the directed tests (SWDEV-212161) were failing when hcc was bumped to 3.0.

[ROCm/hip commit: 6b06911ef1]

projects/hip/include/hip/hcc_detail/device_functions.h

@@ -794,7 +794,7 @@ void *__amdgcn_get_dynamicgroupbaseptr() {
     return __get_dynamicgroupbaseptr();
 }
 
-#if defined(__HCC__) && (__hcc_minor__ < 3)
+#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
 // hip.amdgcn.bc - sync threads
 #define __CLK_LOCAL_MEM_FENCE    0x01
 typedef unsigned __cl_mem_fence_flags;

projects/hip/include/hip/hcc_detail/device_library_decls.h

@@ -80,7 +80,7 @@ extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(vo
 __device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; }
 #endif //__HIP_DEVICE_COMPILE__
 
-#if defined(__HCC__) && (__hcc_minor__ < 3)
+#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
 // __llvm_fence* functions from device-libs/irif/src/fence.ll
 extern "C" __device__ void __llvm_fence_acq_sg(void);
 extern "C" __device__ void __llvm_fence_acq_wg(void);

projects/hip/src/hip_stream.cpp

@@ -31,7 +31,7 @@ THE SOFTWARE.
 //-------------------------------------------------------------------------------------------------
 // Stream
 //
-#if defined(__HCC__) && (__hcc_minor__ < 3)
+#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
 enum queue_priority
 {
     priority_high = 0,
@@ -73,7 +73,7 @@ hipError_t ihipStreamCreate(TlsData *tls, hipStream_t* stream, unsigned int flag
                 // Obtain mutex access to the device critical data, release by destructor
                 LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData());
 
-#if defined(__HCC__) && (__hcc_minor__ < 3)
+#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
                 auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
 #else
                 auto istream = new ihipStream_t(ctx, acc.create_view(Kalmar::execute_any_order, Kalmar::queuing_mode_automatic, (Kalmar::queue_priority)priority), flags);
@@ -242,7 +242,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
     } else if (stream == hipStreamNull) {
         return ihipLogStatus(hipErrorInvalidResourceHandle);
     } else {
-#if defined(__HCC__) && (__hcc_minor__ < 3)
+#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
         *priority = 0;
 #else
         LockedAccessor_StreamCrit_t crit(stream->criticalData());