# CUDAnative: register host memory for pinned memory access

**URL:** https://discourse.julialang.org/t/cudanative-register-host-memory-for-pinned-memory-access/23046
**Category:** GPU
**Tags:** question
**Created:** [April 11, 2019, 5:02pm UTC](https://discourse.julialang.org/t/cudanative-register-host-memory-for-pinned-memory-access/23046 "2019-04-11T17:02:32Z")
**Posts on this page:** 1
**Showing post:** 1

<div class="post-metadata">

### Author: ![samo](https://sea2.discourse-cdn.com/julialang/user_avatar/discourse.julialang.org/samo/32/35398_2.png) [@samo](https://discourse.julialang.org/u/samo)
#### Post date: [April 11, 2019, 5:02pm UTC](https://discourse.julialang.org/t/cudanative-register-host-memory-for-pinned-memory-access/23046/1 "2019-04-11T17:02:32Z")

</div>

Hi all,

how do I register host memory with CUDAnative in order to **enable pinned memory access directly from a GPU kernel**?

In other words, I wonder how in Julia one can do the equivalent of the following CUDA code:

```julia
    float *hbuf_d, *hbuf_h;
    int nx = 512*512*1024;

    // Allocate a host buffer, register it, get a device pointer from it.
    hbuf_h = (float*)malloc((size_t)nx*sizeof(float));
    cudaHostRegister((float*)hbuf_h, (size_t)nx*sizeof(float), cudaHostRegisterMapped);
    cudaHostGetDevicePointer((float**)&hbuf_d, (float*)hbuf_h, 0);

```

This is a snippet from a CUDA code to test the sustained performance of pinned memory access from a GPU kernel (the results are very similar to those of the NVIDIA ‘bandwidthTest’); I add it at the end of this post [1].

Thank you very much!

Sam

[1] CUDA code to test the sustained performance of pinned memory access from a GPU kernel:

```julia
#include <sys/time.h>
#include <stdio.h>

double get_time(){
    struct timeval tp;
    gettimeofday(&tp,NULL);
    return ( (double) tp.tv_sec + (double) tp.tv_usec * 1.e-6 );
}

void cuda_finalize(){
    cudaError_t ce = cudaGetLastError();  
    if(ce != cudaSuccess){ 
        printf("ERROR ON GPU: %s\n", cudaGetErrorString(ce));
    }
}

#define DAT double
//#define DAT float

__global__ void copy_h2d(DAT* hbuf, DAT* dbuf){
    int ix = blockIdx.x*blockDim.x + threadIdx.x;
    dbuf[ix] = hbuf[ix];
}

__global__ void copy_d2h(DAT* hbuf, DAT* dbuf){
    int ix = blockIdx.x*blockDim.x + threadIdx.x;
    hbuf[ix] = dbuf[ix];
}

int main(int argc, char *argv[]){
    int warmup = 3;
    int nx = 512*512*1024; // 512*512*1024*8B = 2GB
    int nt = 10; // 10*2GB = 20GB (total amount of data transfer)
    double t0, time_s, GBs;
    dim3 block, grid;
    DAT *hbuf_d, *hbuf_h, *dbuf_d;
    block.x = 1024; grid.x = nx/block.x;
    block.y = 1; grid.y = 1; 
    block.z = 1; grid.z = 1;

    cudaDeviceReset();

    // Allocate host buffer, register it, get a device pointer from it.
    hbuf_h = (DAT*)malloc((size_t)nx*sizeof(DAT));
    cudaHostRegister((DAT*)hbuf_h, (size_t)nx*sizeof(DAT), cudaHostRegisterMapped); 
    cudaHostGetDevicePointer((DAT**)&hbuf_d, (DAT*)hbuf_h, 0);

    // Allocate device buffer.
    cudaMalloc(&dbuf_d,(size_t)nx*sizeof(DAT));

    // Initialize the host buffer.
    for(int ix=0; ix<nx; ix++){
        hbuf_h[ix] = 1.0;
    }

    // Copy from host to device.
    for (int it=0; it<nt+warmup; it++){
        if (it==warmup) t0 = get_time();
        copy_h2d<<<grid,block>>>(hbuf_d, dbuf_d); cudaDeviceSynchronize();
    }
    time_s = get_time() - t0;
    GBs = 1.0/1024/1024/1024*nt*nx*sizeof(DAT)/time_s;
    printf("h2d: time: %.4f GB/s: %.4f\n", time_s, GBs);

    // Copy from device to host.
    for (int it=0; it<nt+warmup; it++){
        if (it==warmup) t0 = get_time();
        copy_d2h<<<grid,block>>>(hbuf_d, dbuf_d); cudaDeviceSynchronize();
    }
    time_s = get_time() - t0;
    GBs = 1.0/1024/1024/1024*nt*nx*sizeof(DAT)/time_s;
    printf("d2h: time: %.4f; GB/s: %.4f\n", time_s, GBs);

    free(hbuf_h); cudaFree(dbuf_d);
}

```

---

_[View the full topic](https://discourse.julialang.org/t/cudanative-register-host-memory-for-pinned-memory-access/23046)._