memory fragmentation (perhaps?)

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

edit: my /previous/ issue, not my current one was memory corruption. I actually think I got rid of that…that was a different error message entirely. Something along the lines of glibc detected such-and-such* !prev, and then a map of allocated addresses and a stacktrace. Sorry about the confusing title.

I am researching quantum chemistry and plasma physics using a robust, perturbative, renormalized QFT formulated using MaxEnt, Spacetime Algebra/Bivector General Relativity, and active stochastic processes (see e.g. doi:10.1016/j.physletb.2004.03.036, doi:10.1016/j.physletb.2007.10.060, also check out publications by D. Hestenes and E.T. Jaynes). The result will be open-sourced as soon as I can debug it, so feel free to ask for code snippets, etc.

Right now, I’ve got a nasty problem. Using the runtime API, I need to allocate 2 float arrays, each 1286464, asynchronously on a 512mb version of the 8600GT (yes, I know, I’m cheap). Then I need to call my initialization kernel on an input file stored as a uchar*, which has only about 26 registers—I launch 2 streams, each 2 blocks of size (2,2,2). These sizes are subject to as much change as they need given the hardware. I do, however, want to test support for asynchronous transfers and launches if I can, because I made an object-oriented wrapper and am going to write template to abstract the initialization kernel (right now it’s just a dummy that acts as “cudaMemset”). Ideally, more powerful hardware would allow applications to perform multiple simulations without interference.

I don’t know exactly what needs to be fixed. Here is what seems to be the problem code, lines 36-76:

__host__ cystGPUMsim::cystGPUMsim(struct simParams *params, const char* buff){	dp=(cudaPitchedPtr*) malloc(sizeof(cudaPitchedPtr)*(nstreams));	width=params->width;	height=params->height;	depth=params->depth;	dT=params->dT;	dx=params->dx;	n=params->n;	grid=params->grid;	block=params->block;	max=params->max;	nstreams=params->nstreams;	ca_extent.width = width/nstreams;	ca_extent.height = height;	ca_extent.depth = depth;	unsigned char* dbuff;	cudaMalloc((void**) &dbuff, sizeof(buff));	cudaMemcpy(dbuff, buff, sizeof(buff), cudaMemcpyHostToDevice);	stream=(cudaStream_t*) malloc(sizeof(cudaStream_t)*nstreams);	for (int s = 0; s < nstreams; ++s){	dp[s]=make_cudaPitchedPtr( NULL, width/nstreams*sizeof(float), height, depth*4 );	cudaMalloc3D(&(dp[s]), ca_extent);	}	for(int i=0; i<int(width/nstreams); i+=grid.x*block.x){	for(int j=0; j<int(height); i+=grid.y*block.y){	for(int k=0; k<int(depth*4); i+=grid.z*block.z){	for (int s = 0; s < nstreams; ++s){	initialize<<<grid, block, 0, stream[s]>>>((float***) dp[s].ptr, dbuff, i,j,k);	}	cudaThreadSynchronize();	checkCUDAError(" ");	}	}	}	cudaFree(dbuff);	}

Compile-time yields no relevant warnings. What I see on runtime is the following error:

and the gdb stack says:

I would use deviceemu to debug. But right now, if I compile it with device emulation, I receive the following warnings:

Then, when I run it, this odd segmentation fault occurs:

I think this has something to do with my host system’s being single-threaded; I am pretty sure the CUDA Programming Guide said you need at least enough threads on the host to run the device code.

I am running Ubuntu 8.04 32-bit Desktop Ed. with the Cuda 2.1 Beta Toolkit and SDK (my installation doesn’t seem to cause any problems), drivers 180.06 beta, and all of the most recent package updates (=> up-to-date GCC, etc.). I have a Pentium D 805, 1GB of RAM (DDR2, don’t know what speed), and a BFG Tech GeForce 8600 GT OC w/ 512 GB GDDR3 (quite a mouthful). Only nvcc settings are -G, -g, and -arch SM_11, which is the correct target (I checked. I had an atomicAdd that I got rid of, if you must know).

Any thoughts? Appreciate your comments.

2

Can you post full source?

3

sure, sure. here goes:

cyst.h:

/* Copyright (c) 2009 Kevin Daley	The following code is part of cYst.	cYst is free software: you can redistribute it and/or modify	it under the terms of the GNU General Public License as published by	the Free Software Foundation, either version 3 of the License, or	(at your option) any later version.	cYst is distributed in the hope that it will be useful,	but WITHOUT ANY WARRANTY; without even the implied warranty of	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the	GNU General Public License for more details.	You should have received a copy of the GNU General Public License	along with cYst. If not, see <http://www.gnu.org/licenses/>. */ #include <cuda_gl_interop.h> #include <cstdio> #include <cstdlib> #include <cmath> __global__ void initialize(float*** hd, unsigned char* file, int i, int j, int k); extern int resx, resy; namespace cyst{ struct simParams{	float dT;	float n;	float4 dx;	dim3 grid;	dim3 block;	float max;	unsigned int width, height, depth;	int nstreams; }; class cystGPUMsim{ protected:	cudaPitchedPtr *dp;	float dT;	float4 dx;	float n;	dim3 grid;	dim3 block;	cudaStream_t *stream;	int nstreams;	float max;	unsigned int width, height, depth;	cudaExtent ca_extent; public:	__host__ explicit cystGPUMsim(struct simParams *params, const char* buff);	__host__ virtual ~cystGPUMsim();	__host__ void callGPUsim(int frames);	__host__ float* retrieveGPUsim(); }; }

cyst.cu

/* Copyright (c) 2009 Kevin Daley	The following code is part of cYst.	cYst is free software: you can redistribute it and/or modify	it under the terms of the GNU General Public License as published by	the Free Software Foundation, either version 3 of the License, or	(at your option) any later version.	cYst is distributed in the hope that it will be useful,	but WITHOUT ANY WARRANTY; without even the implied warranty of	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the	GNU General Public License for more details.	You should have received a copy of the GNU General Public License	along with cYst. If not, see <http://www.gnu.org/licenses/>. */ #include "cyst.cu" __global__ void initialize(float*** hd, unsigned char* file, int i, int j, int k){	uint3 pg=make_uint3(threadIdx.x+blockIdx.x*blockDim.x+i,threadIdx.x+blockIdx .y*blockDim.y+j,threadIdx.z+blockIdx.z*blockDim.z+k);	hd[pg.x][pg.y][pg.z]=3; } void demo(){ cudaSetDevice(0); struct simParams *parms= new simParams; parms->dT=10; parms->n=.5; parms->dx=make_float4(0.1,0.1,0.1,0.1); dim3 grid (2,1); dim3 block (2,2,2); parms->width=128; parms->height=128; parms->depth=64; parms->max=50.0f; parms->nstreams=2; cudaSetDevice(0);	parms->grid=grid;	parms->block=block;	cyst::cystGPUMsim* sim=new cyst::cystGPUMsim(parms, "k");	sim->callGPUsim(1);	float* H=sim->retrieveGPUsim();	delete sim;	FILE* f=fopen("out.dat","rw+");	for(int i=0; i<256; i++){	for(int j=0; j<256; j++){	for(int k=0; k<256; k++){	fprintf(f, "%f %f %f %f %f %f %f %f\n", i*parms->dx.x, j*parms->dx.y,k*parms->dx.z, k*parms->dx.w, H[(i+j*256+k*256*1024)+3], H[(i+j*256+k*256*1024)], H[(i+j*256+k*256*1024)+1], H[(i+j*256+k*256*1024)+2]);	}}}	fclose(f);	exit(0); }

and demo.cu

/* Copyright (c) 2009 Kevin Daley	The following code is part of cYst.	cYst is free software: you can redistribute it and/or modify	it under the terms of the GNU General Public License as published by	the Free Software Foundation, either version 3 of the License, or	(at your option) any later version.	cYst is distributed in the hope that it will be useful,	but WITHOUT ANY WARRANTY; without even the implied warranty of	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the	GNU General Public License for more details.	You should have received a copy of the GNU General Public License	along with cYst. If not, see <http://www.gnu.org/licenses/>. */ #include "cyst.cu" __global__ void initialize(float*** hd, unsigned char* file, int i, int j, int k){	uint3 pg=make_uint3(threadIdx.x+blockIdx.x*blockDim.x+i,threadIdx.x+blockIdx .y*blockDim.y+j,threadIdx.z+blockIdx.z*blockDim.z+k);	hd[pg.x][pg.y][pg.z]=3; } void demo(){ cudaSetDevice(0); struct simParams *parms= new simParams; parms->dT=10; parms->n=.5; parms->dx=make_float4(0.1,0.1,0.1,0.1); dim3 grid (2,1); dim3 block (2,2,2); parms->width=128; parms->height=128; parms->depth=64; parms->max=50.0f; parms->nstreams=2; cudaSetDevice(0);	parms->grid=grid;	parms->block=block;	cyst::cystGPUMsim* sim=new cyst::cystGPUMsim(parms, "k");	sim->callGPUsim(1);	float* H=sim->retrieveGPUsim();	delete sim;	FILE* f=fopen("out.dat","rw+");	for(int i=0; i<256; i++){	for(int j=0; j<256; j++){	for(int k=0; k<256; k++){	fprintf(f, "%f %f %f %f %f %f %f %f\n", i*parms->dx.x, j*parms->dx.y,k*parms->dx.z, k*parms->dx.w, H[(i+j*256+k*256*1024)+3], H[(i+j*256+k*256*1024)], H[(i+j*256+k*256*1024)+1], H[(i+j*256+k*256*1024)+2]);	}}}	fclose(f);	exit(0); }

and kernel.cu:

/* Copyright (c) 2009 Kevin Daley	The following code is part of cYst.	cYst is free software: you can redistribute it and/or modify	it under the terms of the GNU General Public License as published by	the Free Software Foundation, either version 3 of the License, or	(at your option) any later version.	cYst is distributed in the hope that it will be useful,	but WITHOUT ANY WARRANTY; without even the implied warranty of	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the	GNU General Public License for more details.	You should have received a copy of the GNU General Public License	along with cYst. If not, see <http://www.gnu.org/licenses/>. */ //planck units #define const_pi 3.1415927f #define kC 1.0f #define kB 1.0f #define lightspeed 1.0f #define finestructure 1.0f #include <sm_11_atomic_functions.h> __global__ void kernelGPU(float ***H, float dT, float4 dx, float n, float max, int i, int j, int k){	float mg=0;	uint3 pl=make_uint3(threadIdx.x, threadIdx.y, threadIdx.z);	uint3 pg=make_uint3(threadIdx.x+blockIdx.x*blockDim.x+i,threadIdx.x+blockIdx .y*blockDim.y+j,threadIdx.z+blockIdx.z*blockDim.z+k);	float x=0.0f, y=0.0f, z=0.0f;	__shared__ float shared[2][2][2], E[2][2][2], E2[2][2][2];	E[pl.x][pl.y][pl.z]=H[pg.x][pg.y][pg.z];	E2[pl.x][pl.y][pl.z]=0.0f;	for(float i=0; i<max && i<E[pl.x][pl.y][pl.z]; i+=n){	if((threadIdx.z+1)%4==0){	shared[pl.x][pl.y][pl.z]=i;	mg=sqrtf(pow(shared[pl.x][pl.y][pl.z-3],2.0f)+pow(shared[pl.x][pl.y][pl.z-2],2.0f)+pow(shared[pl.x][pl.y][pl.z-1],2.0f));	x+=shared[pl.x][pl.y][pl.z]*(cosf(mg)-1/mg*shared[pl.x][pl.y][pl.z]*sinf(mg));	}	else if((threadIdx.z+1)%3==0){	shared[pl.x][pl.y][pl.z]=i;	mg=sqrtf(pow(shared[pl.x][pl.y][pl.z+1],2.0f)+pow(shared[pl.x][pl.y][pl.z-2],2.0f)+pow(shared[pl.x][pl.y][pl.z-1],2.0f));	x+=shared[pl.x][pl.y][pl.z]*(cosf(mg)-1/mg*shared[pl.x][pl.y][pl.z]*sinf(mg));	}	else if((threadIdx.z+1)%2==0){	shared[pl.x][pl.y][pl.z]=i;	mg=sqrtf(pow(shared[pl.x][pl.y][pl.z+1],2.0f)+pow(shared[pl.x][pl.y][pl.z+2],2.0f)+pow(shared[pl.x][pl.y][pl.z-1],2.0f));	x+=shared[pl.x][pl.y][pl.z]*(cosf(mg)-1/mg*shared[pl.x][pl.y][pl.z]*sinf(mg));	}	else{	shared[pl.x][pl.y][pl.z]=i;	mg=sqrtf(pow(shared[pl.x][pl.y][pl.z+1],2.0f)+pow(shared[pl.x][pl.y][pl.z+2],2.0f)+pow(shared[pl.x][pl.y][pl.z+3],2.0f));	x+=shared[pl.x][pl.y][pl.z]*(cosf(mg)-1/mg*shared[pl.x][pl.y][pl.z]*sinf(mg));	}	shared[pl.x][pl.y][pl.z]=x;	__syncthreads();	__syncthreads();	if((threadIdx.z+1)%4==0){	z=shared[pl.x][pl.y][pl.z-1];	y=shared[pl.x][pl.y][pl.z-2];	x=shared[pl.x][pl.y][pl.z-3];	}	__syncthreads();	if((threadIdx.z+1)%4==0 && int(fdividef(x,dx.x))>=0 && int(fdividef(x,dx.x))){	E2[pl.x+int(fdividef(x,dx.x))][pl.y+int(fdividef(y,dx.y))][pl.z+4*int(fdividef(z,dx.z))]+=shared[pl.x][pl.y][pl.z];	}	else if((threadIdx.z+1)%4==0){	H[int(fdividef(x,dx.x))+pg.x][int(fdividef(y,dx.y))+pg.y][(int(fdividef(z,dx.z))*4+pg.z)]+=i;	}	E2[pl.x][pl.y][pl.z]-=i;	__syncthreads();	}	H[pg.x][pg.y][pg.z]=E[pl.x][pl.y][pl.z]+E2[pl.x][pl.y][pl.z];	}

Finally, main.cpp:

/* Copyright (c) 2009 Kevin Daley	The following code is part of cYst.	cYst is free software: you can redistribute it and/or modify	it under the terms of the GNU General Public License as published by	the Free Software Foundation, either version 3 of the License, or	(at your option) any later version.	cYst is distributed in the hope that it will be useful,	but WITHOUT ANY WARRANTY; without even the implied warranty of	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the	GNU General Public License for more details.	You should have received a copy of the GNU General Public License	along with cYst. If not, see <http://www.gnu.org/licenses/>. */ extern void demo(); int main(int argc, char** argv){	demo(); }

There you go.

4

bump