@@ -256,7 +256,7 @@ void* Connection::buf_alloc([[maybe_unused]] struct ibv_pd* pd,
|
||||
if (use_gpu_mem) {
|
||||
void* dev_ptr;
|
||||
if (coherent_cq == 1) {
|
||||
#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP
|
||||
#if defined USE_COHERENT_HEAP
|
||||
CHECK_HIP(hipMalloc(reinterpret_cast<void**>(&dev_ptr), size));
|
||||
#else
|
||||
#ifdef HIP_SUPPORTS_MALLOC_UNCACHED
|
||||
|
||||
@@ -40,7 +40,7 @@ namespace rocshmem {
|
||||
|
||||
#if defined USE_MANAGED_HEAP
|
||||
using HEAP_T = HeapMemory<HIPAllocatorManaged>;
|
||||
#elif defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP
|
||||
#elif defined USE_COHERENT_HEAP
|
||||
using HEAP_T = HeapMemory<HIPAllocator>;
|
||||
#elif defined USE_HOST_HEAP
|
||||
using HEAP_T = HeapMemory<HostAllocator>;
|
||||
|
||||
@@ -126,7 +126,7 @@ extern const int gpu_clock_freq_mhz;
|
||||
__device__ __forceinline__ void __roc_inv() { asm volatile("buffer_wbinvl1;"); }
|
||||
|
||||
__device__ __forceinline__ void __roc_flush() {
|
||||
#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP
|
||||
#if defined USE_COHERENT_HEAP
|
||||
#if __gfx90a__
|
||||
asm volatile("s_dcache_wb;");
|
||||
asm volatile("buffer_wbl2;");
|
||||
|
||||
Fai riferimento in un nuovo problema
Block a user