#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
__global__
void kernel2(volatile int * ptr_i)
{ int i=ptr_i[0];
while(i!=1100)
{
if(ptr_i[0]==0)
{
i++;
ptr_i[0]=i;
}
}
syncthreads();
}
int main()
{
int *a; // Pinned memory allocated on the CPU
int *d_a;
int nelem;
unsigned int flags;
size_t bytes;
nelem = 10;
bytes = nelem*sizeof(int);
cudaSetDeviceFlags(cudaDeviceMapHost);
flags = cudaHostAllocMapped;
cudaHostAlloc((void **)&a, bytes, flags);
cudaHostGetDevicePointer((void **)&d_a, (void *)a, 0);
a[0]=0;
kernel2<<<1,1024>>>(d_a);
while(a[0]!=1100)
{
if(a[0]!=0)
{printf("=> %i \n ",a[0]);
a[0]=0;
}
}
}
I modified some things: 1024 threads (the max available according to the programming guide) and less data… What I obtained is in the attached file (don’t care of the .cu extension => needed to upload the file :D)
Check Appendix B.5 of the Programming Guide: __threadfence_system() is needed to make memory changes in mapped host memory visible to the host.
Don’t depend on it too much though: Tim Murray of Nvidia says, it’s not really reliable as there are too many third-party components involved like mainboard chipset and BIOS to get it to work everywhere.
Check Appendix B.5 of the Programming Guide: __threadfence_system() is needed to make memory changes in mapped host memory visible to the host.
Don’t depend on it too much though: Tim Murray of Nvidia says, it’s not really reliable as there are too many third-party components involved like mainboard chipset and BIOS to get it to work everywhere.
I spent a bit of time with it. I got it to work perfectly reliably on my test machines, but there are so many OS/chipset combinations that I am wary of recommending it to others. It’s not something that’s officially supported (in fact it’s expressly forbidden by the documentation), we don’t test it, and that means it could break.
I spent a bit of time with it. I got it to work perfectly reliably on my test machines, but there are so many OS/chipset combinations that I am wary of recommending it to others. It’s not something that’s officially supported (in fact it’s expressly forbidden by the documentation), we don’t test it, and that means it could break.