AI/rocm-systems
2
0
Atdalīts 0
Kods Problēmas Izmaiņu pieprasījumi Darbības Pakotnes Projekti Laidieni Vikivietne Aktivitāte

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

Pārlūkot izejas kodu 
Handled the HCC version check appropriately as few of the directed tests (SWDEV-212161) were failing when hcc was bumped to 3.0.
Šī revīzija ir iekļauta:
satyanveshd
2019-11-20 21:37:52 +05:30
revīziju iesūtīja Maneesh Gupta
vecāks 022ac3cb0a
revīzija d4dde7a27d
3 mainīti faili ar 5 papildinājumiem un 5 dzēšanām
Lejupielādēt ielāpa failu Lejupielādēt izmaiņu failu Izvērst visus failus Savērst visus failus
hipamd/include/hip/hcc_detail/device_functions.h
+1 -1
Parādīt failu
@@ -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;
hipamd/include/hip/hcc_detail/device_library_decls.h
+1 -1
Parādīt failu
@@ -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);
hipamd/src/hip_stream.cpp
+3 -3
Parādīt failu
@@ -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());