atomicAdd causes an illegal memory access

Accelerated Computing CUDA CUDA Programming and Performance
1

The atomicadd in this code is giving me an illegal memory access error, but I don’t know why. Is my usage of atomicAdd incorrect? Even if I comment out the for loop in potential() and simply have it repeatedly perform an atomicAdd on U of 0.0, I still get an error.

#define REALLY_PRINT_TIME
//#define PRINTTIMECLOCK
#include "mycuda.h"
#include <iostream>
#include <string>
#include <stdlib.h>
#include <stdio.h>
#include <math.h>

using namespace std;

const int n=17974;
__device__ __managed__  float data[n][3];


__device__ float distance( float x1, float y1, float z1, float x2, float y2, float z2) {
	float squaresum=(x2-x1)*(x2-x1)+(y2-y1)*(y2-y1)+(z2-z1)*(z2-z1);
	return sqrt(squaresum);
}

__global__ void potential(float *U) {
	const int i= blockIdx.x;
	printf("i==%d\n", i);
	if(i<n) {
		float isum=0.0;
		for(unsigned int j=0; j<n; j++) {
			if(data[i][0]==data[j][0]&&
			   data[i][1]==data[j][1]&&
			   data[i][2]==data[j][2]) {
				continue;
			}
			isum+=distance(data[i][0],data[i][1],data[i][2],
				       data[j][0],data[j][1],data[j][2]);
		}
		atomicAdd(U,1.0/isum);
	}
	//printf("Pot is %.2f\n",*U);
}

void cudamul2() {
	float U=0.0;
	potential<<<n,1>>>(&U);
	cudaError_t cudaerr=cudaDeviceSynchronize();
	if(cudaerr != cudaSuccess) cout << "CUDA Error! "<< cudaGetErrorString(cudaerr) << endl;
	cout << "Potential Energy is " << U << endl;
}
2

You’re passing in a host memory pointer to the host’s U variable, not a device pointer. You need to allocate device memory and pass a pointer to that as your kernel argument, then copy it back to the host to print it.

3

Can I accomplish this using cudaMallocManaged?

4

Got it, thanks!