gistfile1.txt

#include <unistd.h>
#include <sys/time.h>

#include <iostream>
#include <chrono>

#include <omp.h>
#include <stdio.h>
#include <stdlib.h>

#include <hip/hip_runtime.h>

#define BLK_DIM 512


template<typename T>
__device__
static T read(T const* a, uint32_t i) {
  const uint8_t* a_ = (const uint8_t*) a;
  // keeping byte offset in 32 bit unsigned
  // makes it possible to use SGPR-pair + VGPR offset load instruction
  uint32_t byte_off = i * sizeof(T);
  return *((const T*) (a_ + byte_off));
}

template<typename T>
__device__
static void write(T * a, uint32_t i, T val) {
  uint8_t* a_ = (uint8_t*) a;
  // keeping byte offset in 32 bit unsigned
  // makes it possible to use SGPR-pair + VGPR offset store instruction
  uint32_t byte_off = i * sizeof(T);
  *((T*) (a_ + byte_off)) = val;
}




__global__
void my_memcpy(double * __restrict__ dst,  const double *  __restrict__ src, const size_t N){

#if 0
    size_t blockStart = (blockIdx.x * blockDim.x) ;
    double * __restrict__ dst_blk_ptr = &dst[blockStart];
    const double * __restrict__ src_blk_ptr = &src[blockStart]; 
    uint32_t subOffset = threadIdx.x;

    if (blockStart + subOffset < N)
    write(dst_blk_ptr,subOffset, read(src_blk_ptr,subOffset) );


#else
    const size_t i_start = threadIdx.x + blockIdx.x*blockDim.x;
    const size_t i_shift = blockDim.x*gridDim.x;

    for (size_t i = i_start; i < N; i+=i_shift)
       dst[i] = src[i];
#endif

}


int main(){

  double * X, * Y, *Z;
  size_t N = (size_t) 1024*1024*1024/sizeof(double);
  double GB = (double) N*sizeof(double) / (1024.0*1024.0*1024.0);

  int nGPUs = 0;
  hipGetDeviceCount(&nGPUs);
  printf("nGPUs = %d array_size = %g [GB] \n",nGPUs,GB);

  hipSetDevice(0);
  //X = (double*)   omp_target_alloc(N*sizeof(double), omp_get_default_device()) ;
  //X = new double[N]; //as bad as 16 byte alignment in posix_memalign
  //Y = new double[N]; //as bad as 16 byte alignment in posix_memalign

#if defined(USE_SYSTEM_MEMORY)
  posix_memalign((void**) &X, 256, sizeof(double)*N);
  posix_memalign((void**) &Y, 256, sizeof(double)*N);
  posix_memalign((void**) &Z, 256, sizeof(double)*N);
#elif defined(USE_MANAGED_MEMORY)  
  hipMallocManaged((void**) &X, sizeof(double)*N);
  hipMallocManaged((void**) &Y, sizeof(double)*N);
  hipMallocManaged((void**) &Z, sizeof(double)*N);
#else  
  hipMalloc((void**) &X, sizeof(double)*N);
  hipMalloc((void**) &Y, sizeof(double)*N);  
  hipMalloc((void**) &Z, sizeof(double)*N);
#endif

  const int nthreads = BLK_DIM; 
  //int nteams = 440*(1024/nthreads)/2;
  int nteams = (N+nthreads-1)/nthreads;
  int iter  = 4;
  double t1,t2;

 #if 1 //if set to "0"   perfoamnce with "no wait" drops a lot !
  //fill data on CPU
  for (int i = 0; i < iter; ++i){
    t1=omp_get_wtime();
    auto t_start = std::chrono::high_resolution_clock::now();

    #pragma omp parallel for 
    for (size_t i = 0; i < N; ++i){
       X[i] = 0.000001*i;
       Y[i] = 0.000003*i;
       Z[i] = 0.000005*i;
    }
    t2=omp_get_wtime();
    auto t_end = std::chrono::high_resolution_clock::now();
    double seconds =
        std::chrono::duration<double>(t_end - t_start).count();

    printf("CPU memory initialization:  loop time  = %g BW = %g[GB/s]\n",(t2-t1), 3.0*GB/(t2-t1) );
    std::cout << "CPU memory initialization = " << seconds << " seconds, effective BW:  " <<  3.0*GB / seconds  << "[GB/s] \n";

  }

#endif
  //hipMemPrefetchAsync(X,sizeof(double)*N,0);
  //hipMemPrefetchAsync(Y,sizeof(double)*N,0);

  //hipMemAdvise ( (void*) X, sizeof(double)*N, hipMemAdviseSetPreferredLocation, 2 );//seems to have no effect
  //hipMemAdvise ( (void*) Y, sizeof(double)*N, hipMemAdviseSetPreferredLocation, 2 );
  hipDeviceSynchronize();
  
  
  for (int cycle = 0; cycle < 8; ++cycle){
    
    int deviceID = cycle%nGPUs;	  
    hipSetDevice(deviceID);
    printf("using device ID: %d\n",deviceID);

    #ifdef USE_CPU2GPU_PREFETCH
    t1=omp_get_wtime();
    auto t_start = std::chrono::high_resolution_clock::now();

    hipMemPrefetchAsync(X,sizeof(double)*N,deviceID);
    hipMemPrefetchAsync(Y,sizeof(double)*N,deviceID);

    hipDeviceSynchronize();
    //hipStreamSynchronize(0);
    t2=omp_get_wtime();
    auto t_end = std::chrono::high_resolution_clock::now();
    double seconds =
        std::chrono::duration<double>(t_end - t_start).count();

    printf("hipMemPrefetchAsync:   time  = %g BW = %g[GB/s]\n",(t2-t1), 2.0*GB/(t2-t1) );
    #endif

 
  for (int i = 0; i < iter; ++i){
    t1=omp_get_wtime();
    auto t_start = std::chrono::high_resolution_clock::now();

    hipLaunchKernelGGL(my_memcpy,nteams,nthreads,0,0,Y,X,N);
    hipDeviceSynchronize();
    t2=omp_get_wtime();
    auto t_end = std::chrono::high_resolution_clock::now();
    double seconds =
        std::chrono::duration<double>(t_end - t_start).count();

    printf("GPU:  loop time  = %g BW = %g[GB/s]\n",(t2-t1), 2.0*GB/(t2-t1) );
    std::cout << "GPU memory copy time = " << seconds << " seconds, effective BW:  " <<  2.0*GB / seconds  << "[GB/s] \n";

  }

/*
//accuracy check
    #pragma omp parallel for 
    for (size_t i = 0; i < N; ++i){
       if (fabs(0.000001*i-Y[i]) > 1e-10) printf("ERROR X[%zu] = %g, Y[%zu] = %g ",i, 0.000001*i,i,Y[i]);
    }
*/

#if 1
   for (int i = 0; i < iter; ++i){
    t1=omp_get_wtime();
    auto t_start = std::chrono::high_resolution_clock::now();

    #pragma omp parallel for  
    for (size_t i = 0; i < N; ++i)
       Y[i] = X[i];
    t2=omp_get_wtime();
    auto t_end = std::chrono::high_resolution_clock::now();
    double seconds =
        std::chrono::duration<double>(t_end - t_start).count();

    printf("CPU:  loop time  = %g BW = %g[GB/s]\n",(t2-t1), 2.0*GB/(t2-t1) );
    std::cout << "CPU memory copy time = " << seconds << " seconds, effective BW:  " <<  2.0*GB / seconds  << "[GB/s] \n";
  }
#endif


  }

  for (size_t i = N-2; i < N; ++i)
    printf("Y[%zu] = %g\n",i,Y[i]);

  //delete[] DATA;
  

  return 0;
}