user19110 your post is completely off-topic in this thread. Please start your own thread for a completely new topic.
@Robert_Crovella Hello Sir, actually I tried printing values inside as well as outside kernel by writing simple program in cuda for addition of two elements. Below is the attached code.
%%cu
#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <complex.h>
#include <cuComplex.h>
#include <time.h>
__global__ void fc(float fc_x11_1,float fc_x11_2,float *devdx){
devdx[0] = fc_x11_1 + fc_x11_2;
printf("devdx value after addition: %f\n", devdx[i]);
}
int main()
{
printf("hello main\n");
float *dx = (float *)malloc((1)*sizeof(float));
float *devdx;
cudaMalloc((void **)&devdx,(1)*sizeof(float));
float x11[147];
int h;
FILE *fp;
fp = fopen("XYZ1.txt", "r");
for(h=0;h<147;h++)
{
fscanf(fp, "%f", &x11[h]);
}
printf("two nu. %f \t %f",x11[1],x11[5]);
cudaMemcpy(devdx,dx, sizeof(float),cudaMemcpyHostToDevice);
fc<<<1,1>>>(x11[1], x11[5],devdx);
//printf("devdx value : %f\n",devdx);
cudaMemcpy(dx, devdx, sizeof(float),cudaMemcpyDeviceToHost);
printf("dx value : %f\n",dx);
}
But devdx value is not printing correctly. I am running this cuda code on google colab platform.
The code you have posted won’t compile. So it is evidently not the code you are running.
@Robert_Crovella Sir, It was not showing any compilation error on google colab, can you please tell me why it won’t compile…is there any problem with cuda syntax?
You can use a compiler to answer this question. After all, that is what I did. Take the code you have posted (only), and try to compile it.
Apart from the won’t-compile problem, one problem I see in your code is here:
That is host code. Understanding what is wrong there is simply a matter of C/C++ coding, not anything to do with CUDA. If you like, compare it with the way you wrote the printf statement in your kernel code, to see if any of the differences seem important to you. (Careful study of that will also expose the compile error issue.)
My goal here is not about teaching people C/C++ programming.
@Robert_Crovella Yes Sir I know I should not ask C/C++ queries, but actually I ran this code using this below statement also to print the value
printf("dx value : %f\n",*dx);
and this statement to print the address
printf("dx value : %f\n",dx);
Acually that was some issue due to google colab, now that code is running is perfectly fine.
nvcc compiler provided by google colab was not showing any error, that was the main issue.
Thank You
@Robert_Crovella @striker159 Hii, My previous code of 2 elements is running perfectly fine but when I am running the whole loop, then I am getting all the elements as 0 as the output.
Please can you tell me that what according to you might be wrong in my CUDA code for subtraction of all x11 values(2 values at a time) inside my code.
This was the code written by me :
%%cu
#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <complex.h>
#include <cuComplex.h>
#include <time.h>
__global__ void fc(float *x11,float *dx){
int row = blockIdx.y*blockDim.y+threadIdx.y;
int col = blockIdx.x*blockDim.x+threadIdx.x;
if((row < 147) && (col < 147)){
dx[(row*147)+col] = x11[row] - x11[col];
}
}
int main()
{
float *dx = (float*)malloc((147*147)*sizeof(float));
float *x11 = (float*)malloc(147*sizeof(float));
float *devdx, *d_x11;
cudaMalloc(&devdx,(147*147)*sizeof(float));
cudaMalloc(&d_x11,(147)*sizeof(float));
int h;
FILE *fp;
if ((fp = fopen("XYZ1.txt","r")) == NULL){
printf("Error! opening file");
exit(1);
}
for(h=0;h<147;h++){
fscanf(fp, "%f", (x11+h));
}
cudaMemcpy(d_x11, x11, (147*sizeof(float)), cudaMemcpyHostToDevice);
dim3 blocksize(147,147);
dim3 threadsize(1,1);
fc<<<blocksize,threadsize>>>(d_x11,devdx);
cudaMemcpy(dx, devdx,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
int i,j;
for(i=0;i<147;i++){
for(j=0;j<147;j++){
printf("%f\t",dx[i*147+j]);
}
printf("\n");
}
return 0;
}
This was the following output which I was getting:
What happens if you follow the instructions I gave you here ? If no errors are reported then my best guess is your input file (which I don’t have) is all zeroes.
I’ve suggested to you twice now that you should do proper cuda error checking (demonstrate it in any code you post) and also run with one of the sanitizers and report the output of that tool when asking for help.
I probably won’t be able to help you further with questions where you don’t demonstrate those things.
If I run your code with synthetic data, the output looks sensible to me:
$ cat t1931.cu
#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <complex.h>
#include <cuComplex.h>
#include <time.h>
__global__ void fc(float *x11,float *dx){
int row = blockIdx.y*blockDim.y+threadIdx.y;
int col = blockIdx.x*blockDim.x+threadIdx.x;
if((row < 147) && (col < 147)){
dx[(row*147)+col] = x11[row] - x11[col];
}
}
int main()
{
float *dx = (float*)malloc((147*147)*sizeof(float));
float *x11 = (float*)malloc(147*sizeof(float));
float *devdx, *d_x11;
cudaMalloc(&devdx,(147*147)*sizeof(float));
cudaMalloc(&d_x11,(147)*sizeof(float));
for(int h=0;h<147;h++){
x11[h] = h;
}
cudaMemcpy(d_x11, x11, (147*sizeof(float)), cudaMemcpyHostToDevice);
dim3 blocksize(147,147);
dim3 threadsize(1,1);
fc<<<blocksize,threadsize>>>(d_x11,devdx);
cudaMemcpy(dx, devdx,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
int i,j;
for(i=0;i<7;i++){
for(j=0;j<7;j++){
printf("%f\t",dx[i*147+j]);
}
printf("\n");
}
return 0;
}
$ nvcc -o t1931 t1931.cu
$ cuda-memcheck ./t1931
========= CUDA-MEMCHECK
0.000000 -1.000000 -2.000000 -3.000000 -4.000000 -5.000000 -6.000000
1.000000 0.000000 -1.000000 -2.000000 -3.000000 -4.000000 -5.000000
2.000000 1.000000 0.000000 -1.000000 -2.000000 -3.000000 -4.000000
3.000000 2.000000 1.000000 0.000000 -1.000000 -2.000000 -3.000000
4.000000 3.000000 2.000000 1.000000 0.000000 -1.000000 -2.000000
5.000000 4.000000 3.000000 2.000000 1.000000 0.000000 -1.000000
6.000000 5.000000 4.000000 3.000000 2.000000 1.000000 0.000000
========= ERROR SUMMARY: 0 errors
$
Don’t assume this means it should be OK in your setup. If you happen to get a colab instance with a K80 GPU in it, the problem may be that you are not compiling the code correctly. If you had done proper CUDA error checking, that error would be reported when you ran the code. But based on the code you have actually posted here, you wouldn’t get any indication of that sort of problem.
@Robert_Crovella Hi Sir, You have given instructions to check errors, after going through those resources/ documentations, I got some basic idea about how to check errors and I have included those statements inside my code which I have posted below::
%%cu
#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <complex.h>
#include <cuComplex.h>
#include <time.h>
inline cudaError_t checkCudaErr(cudaError_t err, const char* msg) {
if (err != cudaSuccess) {
fprintf(stderr, "CUDA Runtime error at %s: %s\n", msg, cudaGetErrorString(err));
}
return err;
}
__global__ void fc(float *x11,float *dx){
int row = blockIdx.y*blockDim.y+threadIdx.y;
int col = blockIdx.x*blockDim.x+threadIdx.x;
if((row < 147) && (col < 147)){
dx[(row*147)+col] = x11[row] - x11[col];
}
}
int main()
{
float *dx = (float*)malloc((147*147)*sizeof(float));
float *x11 = (float*)malloc(147*sizeof(float));
float *devdx, *d_x11;
cudaMalloc(&devdx,(147*147)*sizeof(float));
cudaMalloc(&d_x11,(147)*sizeof(float));
//float x11[147];
int h;
FILE *fp;
if ((fp = fopen("XYZ1.txt","r")) == NULL){
printf("Error! opening file");
exit(1);
}
for(h=0;h<147;h++){
fscanf(fp, "%f", (x11+h));
}
cudaMemcpy(d_x11, x11, (147*sizeof(float)), cudaMemcpyHostToDevice);
dim3 blocksize(147,147);
dim3 threadsize(1,1);
fc<<<blocksize,threadsize>>>(d_x11,devdx);
checkCudaErr(cudaDeviceSynchronize(), "Syncronization");
checkCudaErr(cudaGetLastError(), "GPU");
cudaError err = cudaMemcpy(dx, devdx,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
if(err!=cudaSuccess){
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err));
}
int i,j;
for(i=0;i<147;i++){
for(j=0;j<147;j++){
printf("%f\t",dx[i*147+j]);
}
printf("\n");
}
return 0;
}
Finally I am getting correct ouput values.
Is this the correct way of checking errors? I wanted to study more about cuda error checking and debugging so that I can effectively parallelise the big codes involving multiple for loops.
If I were doing it, I would check every single CUDA runtime API call for an error. You are not doing that.
If I were doing it, I would follow the instructions here (a link I have already posted to you) exactly for error checking after a kernel. You are not doing that.
The example you have here is certainly a lot better than no error checking at all, and it might be sufficient for most things. But I believe there is basically one and only one “proper” way to do rigorous error checking:
- Check every CUDA call.
- Check every kernel call, properly catching launch errors as well as execution errors (that is why there are 2 checks after a kernel call.
For a better treatment, review the material here. (which I have already also suggested.)
I suggest studying carefully the resources I have now suggested to you multiple times, before asking further questions about error checking.
@Robert_Crovella Hello Sir, As per the suggestions and resources provided by you in this thread, I have gone through the stack overflow link. I have incorporated those statements for checking errors related to kernel launch, cudaMemcpy , etc.
As for using cuda-gdb , hardware access is required, so not able to use cuda-gdb as I am working on colab instance of gpu.
Below is the code which I have written in cuda for parallelisation :
%%cu
#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <complex.h>
#include <cuComplex.h>
#include <time.h>
inline cudaError_t checkCudaErr(cudaError_t err, const char* msg) {
if (err != cudaSuccess) {
fprintf(stderr, "CUDA Runtime error at %s: %s\n", msg, cudaGetErrorString(err));
}
return err;
}
__device__ cuDoubleComplex my_complex_exp(cuDoubleComplex ex1){
cuDoubleComplex comp_1_w;
double s, c;
double e = exp(ex1.x);
sincos(ex1.y, &s, &c);
comp_1_w.x = c * e;
comp_1_w.y = s * e;
return comp_1_w;
}
__host__ __device__ void legendre(int lmax, int l, float *theta_dummy, float Lmn[11])
{
int u,i,size_x,jj,hh,ind,nind;
int m_5,j;
float c,tol;
float htol=2.2251e-308;
float inf = 1e+10;
float x,rootn[21],twocot;
float s,P[11],sn;
float tmp1,tmp2,tmp3;
float trm1, trm2, trm3, trm4, trm5;
float t,t11;
for (i=0;i<11;i++)
{
Lmn[i]=0.0;
}
x=cosf((*theta_dummy));
float vec_u[21] = {0.0,1.0,2.0,3.0,4.0,5.0,6.0,7.0,8.0,9.0,10.0,11.0,12.0,13.0,14.0,15.0,16.0,17.0,18.0,19.0,20.0};
for (u = 0; u<21; u++)
{
rootn[u] = 0.0;
}
for (u = 0; u<19; u++)
{
if (u<2*l+1){
rootn[u] = sqrtf(vec_u[u]);
}
}
s = sqrtf(1-(x*x));
for (i = 0; i<9; i++)
{
if (i<l){
P[i] = 0.0;
}
}
twocot = x;
if(x == -1)
{
twocot = inf;
}
else if(x == 1)
{
twocot = -inf;
}
else if(x != 0)
{
twocot = -2*x/s;
}
sn = 0.0;
float xyz=1.0;
int sss;
for(sss=1; sss<=9; sss++)
{
if(sss<=l){
xyz *= (-s);}
}
sn = xyz;
tol = sqrtf(htol);
hh=0;
nind=0;
if ((x != 1) && ((float)(fabsf(sn)) >= tol))
{
nind = hh+1;
hh = hh + 1;
}
if (hh > 0)
{
float d[15] = {0.5,0.25,0.166666667,0.125,0.1,0.083333333,0.071428571,0.0625,0.055555556,0.05,0.045454545,0.041666667,0.038461538,0.035714286,0.033333333};
float dd;
c = 1;
for (i =0; i<9; i++)
{
if (i<l){
//dd = (1-1/d[i]);
dd = (1-d[i]);
c *= dd;
}
}
P[l] = sqrtf(c)*sn; //10
P[l-1] = P[l]*twocot*l/rootn[2*l]; //9
tmp1 = P[l];
tmp2 = P[l-1];
for (m_5 = 7; m_5>=0; m_5+=-1)
{
if (m_5<=l-2)
{
trm1 = tmp2*twocot*(m_5+1);
trm2 = tmp1*rootn[l+m_5+2];
trm3 = rootn[l-m_5-1];
trm4 = rootn[l+m_5+1];
trm5 = rootn[l-m_5];
tmp3 = (trm1-trm2*trm3)/(trm4*trm5);
P[m_5] = tmp3;
tmp1 = tmp2;
tmp2 = tmp3;
}
}
}
for (i=0; i<10; i++)
{
if (i<l+1)
{
t = P[i];
Lmn[i] = t;
}
}
float abc=1.0;
int uuu;
for(uuu=1; uuu<=9; uuu++){
if(uuu<=l){
abc *= x;}
}
if (s == 0)
{
Lmn[0] = abc;
}
for (m_5=1; m_5<=8; m_5++)
{
if (m_5<=(l-1))
{
c = 1;
for (j = 1; j<=16; j++)
{
if (j<=2*m_5)
{
c*=rootn[l-m_5+2+j-2];
}
}
t = Lmn[m_5];
t11 = c*t;
Lmn[m_5] = t11;
}
}
c = 1;
for (i = 0; i<18; i++)
{
if (i<2*l)
{
c *= rootn[i+1];
}
}
Lmn[l] *= c;
}
__global__ void fc(float *x11,float *dx,float *y11,float *dy,float *z11,float *dz,float *r,float *pcolhea1,float *fophi1,float *rcutfun,float *term1,float *G_p,float *func1,cuDoubleComplex *ex1,cuDoubleComplex *comp_1_w,cuDoubleComplex *ex2,cuDoubleComplex *ex3,cuDoubleComplex *mult1,float *cons1,cuDoubleComplex *useY1){
cuDoubleComplex p_ch_p,xt2;
cuDoubleComplex temp1,temp2;
cuDoubleComplex q = make_cuDoubleComplex(0,1);
int l,k2,n,m,b_m,ll,mm,m_1,b_mm,ju,lmax=10,etaang=5,natoms=147;
float Lmn_all[147][10][11],Lmn[11],rcut = 8.0,pi=3.141592653589793,rcut_inv = 0.125;
float Lmmnn;
float netaa[5]={0.0028,0.0040,0.0110,0.0280,0.059};
/*for(int i=0;i<10;i++){
for(int j=0;j<21;j++){
printf("%f\n",cons1[i*10+j]);
}
}*/
int h = blockIdx.y*blockDim.y+threadIdx.y;
int jt = blockIdx.x*blockDim.x+threadIdx.x;
if((h < 147) && (jt < 147) && (h!=jt)){
dx[(h*147)+jt] = x11[h] - x11[jt];
dy[(h*147)+jt] = y11[h] - y11[jt];
dz[(h*147)+jt] = z11[h] - z11[jt];
r[(h*147)+jt] = sqrtf(dx[(h*147)+jt]*dx[(h*147)+jt] + dy[(h*147)+jt]*dy[(h*147)+jt] + dz[(h*147)+jt]*dz[(h*147)+jt]);
pcolhea1[(h*147)+jt] = acosf(dz[(h*147)+jt]/r[(h*147)+jt]);
for (l=0;l<lmax;l++)
{
if (l==0)
{
Lmn[0]=1.0;
Lmn_all[jt][0][0] = 1.0;
//*(Lmn_all + jt*110) = 1.0;
}
else
{
legendre(lmax,l,pcolhea1,Lmn);
}
for(b_m=0;b_m<19;b_m++)
{
if(b_m<2*l+1)
{
m= -l + b_m;
if (m >= 0)
{
Lmn_all[jt][l][m]= Lmn[m];
//*(Lmn_all + jt*110 + l*11 + m) = Lmn[m];
}
}
}
}
k2=0;
for(n=0;n<etaang;n++) // loop 15.2
{
for(ll=0;ll<lmax;ll++)
{
p_ch_p= make_cuDoubleComplex(0,0);
for(b_mm=0;b_mm<19;b_mm++)
{
if(b_mm<2*ll+1)
{
xt2 = make_cuDoubleComplex(0,0);
int ju = blockIdx.x*blockDim.x+threadIdx.x;
if(ju<natoms)
{
if(h != ju)
{
dx[(h*147)+ju]=(x11[h]-x11[ju]);
dy[(h*147)+ju]=(y11[h]-y11[ju]);
dz[(h*147)+ju]=(z11[h]-z11[ju]);
r[(h*147)+ju]=sqrtf(dx[(h*147)+ju]*dx[(h*147)+ju] + dy[(h*147)+ju]*dy[(h*147)+ju] + dz[(h*147)+ju]*dz[(h*147)+ju]);
fophi1[(h*147)+ju]= atan2f(dy[(h*147)+ju],dx[(h*147)+ju]);
if (r[(h*147)+ju] <= rcut)
{
rcutfun[(h*147)+ju] = 0.5*(cosf(pi*r[(h*147)+ju]*rcut_inv)+1);
}
else
{
rcutfun[(h*147)+ju] = 0.0;
}
func1[(h*147)+ju] = expf(-netaa[n]*r[(h*147)+ju]*r[(h*147)+ju]);
term1[(h*147)+ju] = func1[(h*147)+ju]*rcutfun[(h*147)+ju];
G_p[(h*147)+ju] = term1[(h*147)+ju];
mm = -ll+b_mm;
m_1 = abs(mm);
if (mm<0)
{
temp1 = cuCmul(q, make_cuDoubleComplex(-m_1,0));
ex1[(h*147)+ju] = cuCmul(temp1, make_cuDoubleComplex(fophi1[(h*147)+ju],0));
comp_1_w[(h*147)+ju] = my_complex_exp(ex1[(h*147)+ju]);
}
else
{
temp2 = cuCmul(q, make_cuDoubleComplex(m_1,0));
ex2[(h*147)+ju] = cuCmul(temp2, make_cuDoubleComplex(fophi1[(h*147)+ju],0));
ex3[(h*147)+ju] = my_complex_exp (ex2[(h*147)+ju]);
float result = 1.0;
int zz;
for(zz=9; zz>=1; zz--)
{
if(zz<=mm)
{
result *= (-1.0);
}
}
comp_1_w[(h*147)+ju] = cuCmul(ex3[(h*147)+ju], make_cuDoubleComplex(result, 0));
}
mult1[(h*147)+ju] = cuCmul(comp_1_w[(h*147)+ju], make_cuDoubleComplex(cons1[(ll*10)+b_mm], 0));
Lmmnn = Lmn_all[ju][ll][m_1];
//Lmmnn = *(Lmn_all + ju*110 + ll*11 + m_1);
useY1[(h*147)+ju] = cuCmul(mult1[(h*147)+ju], make_cuDoubleComplex(Lmmnn, 0));
}
}
}
}
}
}
}
}
int main()
{
float *dx = (float*)malloc((147*147)*sizeof(float));
float *dy = (float*)malloc((147*147)*sizeof(float));
float *dz = (float*)malloc((147*147)*sizeof(float));
float *x11 = (float*)malloc(147*sizeof(float));
float *y11 = (float*)malloc(147*sizeof(float));
float *z11 = (float*)malloc(147*sizeof(float));
float *r = (float*)malloc((147*147)*sizeof(float));
float *pcolhea1 = (float*)malloc((147*147)*sizeof(float));
float *fophi1 = (float*)malloc((147*147)*sizeof(float));
float *rcutfun = (float*)malloc((147*147)*sizeof(float));
float *term1 = (float*)malloc((147*147)*sizeof(float));
float *G_p = (float*)malloc((147*147)*sizeof(float));
float *func1 = (float*)malloc((147*147)*sizeof(float));
cuDoubleComplex *ex1 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *comp_1_w = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *ex2 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *ex3 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *mult1 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *useY1 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
float *cons11= (float*)malloc(100*sizeof(float));
float *cons1 = (float*)malloc(210*sizeof(float));
FILE * fp1;
fp1 = fopen ("cons11.txt", "r");
int index = 0;
while(fscanf(fp1, "%f",(cons11+index))==1)
index++;
fclose(fp1);
int l,m;
for (l=0;l<10;l++)
{
for(m=0;m<21;m++)
{
if(m<2*l+1)
{
cons1[(l*10)+m] = cons11[(l*l)+m];
}
else
{
cons1[(l*10)+m] = 0.0;
}
}
}
float *devdx, *d_x11, *devdy, *d_y11, *devdz, *d_z11, *d_r, *d_pcolhea1, *d_fophi1, *d_rcutfun, *d_term1, *d_G_p, *d_func1,*d_cons1;
cuDoubleComplex *d_ex1, *d_ex2, *d_ex3, *d_comp_1_w, *d_mult1, *d_useY1;
cudaMalloc(&devdx,(147*147)*sizeof(float));
cudaMalloc(&d_x11,(147)*sizeof(float));
cudaMalloc(&devdy,(147*147)*sizeof(float));
cudaMalloc(&d_y11,(147)*sizeof(float));
cudaMalloc(&devdz,(147*147)*sizeof(float));
cudaMalloc(&d_z11,(147)*sizeof(float));
cudaMalloc(&d_r,(147*147)*sizeof(float));
cudaMalloc(&d_pcolhea1,(147*147)*sizeof(float));
cudaMalloc(&d_fophi1,(147*147)*sizeof(float));
cudaMalloc(&d_rcutfun,(147*147)*sizeof(float));
cudaMalloc(&d_term1,(147*147)*sizeof(float));
cudaMalloc(&d_G_p,(147*147)*sizeof(float));
cudaMalloc(&d_func1,(147*147)*sizeof(float));
cudaMalloc(&d_ex1,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_ex2,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_ex3,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_comp_1_w,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_mult1,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_useY1,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_cons1,(210)*sizeof(float));
//float x11[147];
int h;
FILE *fp;
if ((fp = fopen("XYZ1.txt","r")) == NULL){
printf("Error! opening file");
exit(1);
}
for(h=0;h<147;h++){
fscanf(fp, "%f", (x11+h));
}
for(h=147;h<294;h++)
{
fscanf(fp, "%f", (y11+h-147));
}
for(h=294;h<441;h++)
{
fscanf(fp, "%f", (z11+h-294));
}
cudaMemcpy(d_x11, x11, (147*sizeof(float)), cudaMemcpyHostToDevice);
cudaMemcpy(d_y11, y11, (147*sizeof(float)), cudaMemcpyHostToDevice);
cudaMemcpy(d_z11, z11, (147*sizeof(float)), cudaMemcpyHostToDevice);
cudaMemcpy(d_cons1, cons1, (210*sizeof(float)), cudaMemcpyHostToDevice);
//cudaMemcpy(d_r, r, (147*147*sizeof(float)), cudaMemcpyHostToDevice);
dim3 blocksize(147,147);
dim3 threadsize(1,1);
fc<<<blocksize,threadsize>>>(d_x11,devdx,d_y11,devdy,d_z11,devdz,d_r,d_pcolhea1,d_fophi1,d_rcutfun,d_term1,d_G_p,d_func1,d_ex1,d_comp_1_w,d_ex2,d_ex3,d_mult1,d_cons1,d_useY1);
cudaDeviceSynchronize();
checkCudaErr(cudaDeviceSynchronize(), "Syncronization");
checkCudaErr(cudaGetLastError(), "GPU");
checkCudaErr(cudaPeekAtLastError(), "GPUpeek");
cudaError err1 = cudaMemcpy(dx, devdx,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err2 = cudaMemcpy(dy, devdy,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err3 = cudaMemcpy(dz, devdz,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err4 = cudaMemcpy(r, d_r,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err5 = cudaMemcpy(pcolhea1, d_pcolhea1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err6 = cudaMemcpy(fophi1, d_fophi1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err7 = cudaMemcpy(rcutfun, d_rcutfun,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err8 = cudaMemcpy(term1, d_term1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err9 = cudaMemcpy(G_p, d_G_p,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err10 =cudaMemcpy(func1, d_func1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err11 =cudaMemcpy(ex1, d_ex1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaError err12 =cudaMemcpy(comp_1_w, d_comp_1_w,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(ex2, d_ex2,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(ex3, d_ex3,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
checkCudaErr(cudaMemcpy(mult1, d_mult1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost),"str1");
checkCudaErr(cudaMemcpy(useY1, d_useY1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost),"str2");
//cudaError err13 = cudaMemcpy(mult1, d_mult1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
//cudaError err14 = cudaMemcpy(useY1, d_useY1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
if((err1 || err2 || err3 || err4 || err5 || err6 || err7 || err8 || err9 || err10 || err11 || err12)!=cudaSuccess){
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err1));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err2));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err3));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err4));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err5));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err6));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err7));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err8));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err9));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err10));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err11));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err12));
//printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err13));
//printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err14));
}
int i,j;
for(i=0;i<147;i++){
for(j=0;j<147;j++){
printf("%f %f\n", mult1[(i*147)+j].x, mult1[(i*147)+j].y);
}
}
return 0;
}
After adding those error checking statements , still I am getting the values of mult1 as zeros and not getting any error message.
Also, I am getting the values of mult1 inside device but not inside main. Can you please suggest any possible reason for that.
Thanks a lot in advance
That is not correct.
I cannot help with code that requires me to open files that I don’t have. When I help other people, I usually use the same tools I have suggested to you: sanitizers, debuggers, etc. If I can’t run your code, I can’t use those tools.
Good luck!
@Robert_Crovella
Dear Sir, my code involves only 2 txt files. Please can I mail those two files to you so that you can run this code using your tools.
I don’t want to stop learning after reaching so far. Please help me.
It’s my humble request.
Thank You.
When you need to ask for help on an online forum like this, I think its a good skill to have to learn how to write a problem example that doesn’t depend on external files. That would be my suggestion here.
@Robert_Crovella Sure Sir, I get your point. Please find attached the external files below. I will definitely try to learn how to write problem example that doesn’t depend on external files.
Please help me this time.
when I print mult1 variable inside __global __ kernel, I’m getting non zero values but when I printed mult1 variable inside main() , I’m getting all values as zeros.
cons11.txt (2.3 KB)
XYZ1.txt (8.6 KB)
Your kernel has loops in it. This means that you are setting a given mult1 value in your kernel multiple times. So the very first value you set may be non-zero, or some intermediate values you set may be non-zero, but the final values you set are all printing out (in kernel code) as plus or minus zero.
The things you are printing outside the kernel are the final values, so these match the final values you set in the kernel code. So your observation is not surprising to me.
For a forum, in my opinion, debugging is intractable on such a large code. Furthermore, printout of 147*147 values is not necessary or tractable when you are trying to demonstrate a problem. You could just as easily print out a single member of the mult1 array in both device code and host code to clearly demonstrate the thing you are asking about. I think if you had used in-kernel printf here (which is what I did) the reason for your observation would have already been obvious.
Simplify your code, reduce the scope, and reduce the printout, before asking for help, is my suggestion.
I don’t believe it is related to what you are asking about, but when I compile your code one of the warnings I get is like this:
t1.cu(34): warning: floating-point value does not fit in required floating-point type
I don’t recommend ignoring such a warning. If you don’t know what size of numbers fits in a float variable, that is C++ question you can find answers to elsewhere.
@Robert_Crovella Hello Sir, I have removed that warning as suggested by you. Also, I tried to print only single/ few members of mult1 array in both device and host. It seemed obvious that both values will match and it is matching. But, when I am running the corresponding C code, I am getting different values of mult1 array in comparison to when I am running the code in CUDA . The final values should match in both the cases. Till comp_1_w, all values are matching.
What can be the possible reason sir?
Please give your valuable insights.
For your reference, corresponding C code file is attached below:
tot_cnlm_23nov.c (10.2 KB)
Thank you so much for your time.
@Robert_Crovella Merry Christmas Sir, hope you are fine.
Actually I have written the following code in CUDA but I was getting error as :
CUDA Runtime error at GPU: invalid argument
The Code was as follows:
%%cu
#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <complex.h>
#include <cuComplex.h>
#include <time.h>
inline cudaError_t checkCudaErr(cudaError_t err, const char* msg) {
if (err != cudaSuccess) {
fprintf(stderr, "CUDA Runtime error at %s: %s\n", msg, cudaGetErrorString(err));
}
return err;
}
float elapsed = 0;
cudaEvent_t start, stop;
cudaEvent_t end;
__device__ cuDoubleComplex my_complex_exp(cuDoubleComplex ex1){
cuDoubleComplex comp_1_w;
double s, c;
double e = exp(ex1.x);
sincos(ex1.y, &s, &c);
comp_1_w.x = c * e;
comp_1_w.y = s * e;
return comp_1_w;
}
__host__ __device__ void legendre(int lmax, int l, float dummy, float Lmn[11])
{
int u,i,size_x,jj,hh,ind,nind;
int m_5,j;
float c,tol;
float htol=2.2251e-308;
float inf = 1e+10;
float x,rootn[21],twocot;
float s,P[11],sn;
float tmp1,tmp2,tmp3;
float trm1, trm2, trm3, trm4, trm5;
float t,t11;
for (i=0;i<11;i++)
{
Lmn[i]=0.0;
}
x=cosf((dummy));
//printf("lmax,l,dummy %d\t%d\t%f\n",lmax,l,dummy);
float vec_u[21] = {0.0,1.0,2.0,3.0,4.0,5.0,6.0,7.0,8.0,9.0,10.0,11.0,12.0,13.0,14.0,15.0,16.0,17.0,18.0,19.0,20.0};
for (u = 0; u<21; u++)
{
rootn[u] = 0.0;
}
for (u = 0; u<19; u++)
{
if (u<2*l+1){
rootn[u] = sqrtf(vec_u[u]);}
}
s = sqrtf(1-(x*x));
for (i = 0; i<9; i++)
{
if (i<l){
P[i] = 0.0;}
}
twocot = x;
if(x == -1)
{
twocot = inf;
}
else if(x == 1)
{
twocot = -inf;
}
else if(x != 0)
{
twocot = -2*x/s;
}
sn = 0.0;
float xyz=1.0;
int sss;
for(sss=1; sss<=9; sss++)
{
if(sss<=l){
xyz *= (-s);}
}
sn = xyz;
tol = sqrtf(htol);
hh=0;
nind=0;
if ((x != 1) && ((float)(fabsf(sn)) >= tol))
{
nind = hh+1;
hh = hh + 1;
}
if (hh > 0)
{
float d[15] = {0.5,0.25,0.166666667,0.125,0.1,0.083333333,0.071428571,0.0625,0.055555556,0.05,0.045454545,0.041666667,0.038461538,0.035714286,0.033333333};
float dd;
c = 1;
for (i =0; i<9; i++)
{
if (i<l){
//dd = (1-1/d[i]);
dd = (1-d[i]);
c *= dd;
}
}
P[l] = sqrtf(c)*sn; //10
P[l-1] = P[l]*twocot*l/rootn[2*l]; //9
tmp1 = P[l];
tmp2 = P[l-1];
for (m_5 = 7; m_5>=0; m_5+=-1)
{
if (m_5<=l-2)
{
trm1 = tmp2*twocot*(m_5+1);
trm2 = tmp1*rootn[l+m_5+2];
trm3 = rootn[l-m_5-1];
trm4 = rootn[l+m_5+1];
trm5 = rootn[l-m_5];
tmp3 = (trm1-trm2*trm3)/(trm4*trm5);
P[m_5] = tmp3;
tmp1 = tmp2;
tmp2 = tmp3;
}
}
}
for (i=0; i<10; i++)
{
if (i<l+1)
{
t = P[i];
Lmn[i] = t;
}
}
float abc=1.0;
int uuu;
for(uuu=1; uuu<=9; uuu++){
if(uuu<=l){
abc *= x;}
}
if (s == 0)
{
Lmn[0] = abc;
}
for (m_5=1; m_5<=8; m_5++)
{
if (m_5<=(l-1))
{
c = 1;
for (j = 1; j<=16; j++)
{
if (j<=2*m_5)
{
c*=rootn[l-m_5+2+j-2];
}
}
t = Lmn[m_5];
t11 = c*t;
Lmn[m_5] = t11;
}
}
c = 1;
for (i = 0; i<18; i++)
{
if (i<2*l)
{
c *= rootn[i+1];
}
}
Lmn[l] *= c;
for(i=0;i<11;i++){
//printf("Lmn values : %f\n",Lmn[i]);
}
}
__global__ void fc(float *x11,float *dx,float *y11,float *dy,float *z11,float *dz,float *r,float *pcolhea1,float *fophi1,float *rcutfun,float *term1,float *G_p,float *func1,cuDoubleComplex *ex1,cuDoubleComplex *comp_1_w,cuDoubleComplex *ex2,cuDoubleComplex *ex3,cuDoubleComplex *mult1,float *cons1,cuDoubleComplex *useY1,cuDoubleComplex *mult4,cuDoubleComplex *xt6, cuDoubleComplex *temp3, cuDoubleComplex *p_ch_p){
cuDoubleComplex temp1,temp2,tot_cnlm_sum_p[147][5][10][21],xt2;
cuDoubleComplex q = make_cuDoubleComplex(0,1);
int l,k2,n,m,b_m,ll,mm,m_1,b_mm,ju,lmax=10,etaang=5,natoms=147;
float Lmn_all[147][10][11],Lmn[11],rcut = 8.0,pi=3.141592653589793,rcut_inv = 0.125;
float Lmmnn;
float netaa[5]={0.0028,0.0040,0.0110,0.0280,0.059};
float dummy;
int h = blockIdx.y*blockDim.y+threadIdx.y;
int jt = blockIdx.x*blockDim.x+threadIdx.x;
if((h < 147) && (jt < 147) && (h!=jt)){
dx[(h*147)+jt] = x11[h] - x11[jt];
dy[(h*147)+jt] = y11[h] - y11[jt];
dz[(h*147)+jt] = z11[h] - z11[jt];
r[(h*147)+jt] = sqrtf(dx[(h*147)+jt]*dx[(h*147)+jt] + dy[(h*147)+jt]*dy[(h*147)+jt] + dz[(h*147)+jt]*dz[(h*147)+jt]);
pcolhea1[(h*147)+jt] = acosf(dz[(h*147)+jt]/r[(h*147)+jt]);
for(int i=0;i<147;i++){
for(int j=0;j<147;j++){
if(i!=j){
dummy = pcolhea1[i*147+j];
for (l=0;l<lmax;l++)
{
if (l==0)
{
Lmn[0]=1.0;
Lmn_all[jt][0][0] = 1.0;
//*(Lmn_all + jt*110) = 1.0;
}
else
{
legendre(lmax,l,dummy,Lmn);
}
for(b_m=0;b_m<19;b_m++)
{
if(b_m<2*l+1)
{
m= -l + b_m;
if (m >= 0)
{
Lmn_all[jt][l][m]= Lmn[m];
//printf("Lmn_all : %f\n",Lmn_all[jt][l][m]);
//*(Lmn_all + jt*110 + l*11 + m) = Lmn[m];
}
}
}
}
}
}
}
k2=0;
for(n=0;n<etaang;n++) // loop 15.2
{
for(ll=0;ll<lmax;ll++)
{
p_ch_p[(h*147)+ju] = make_cuDoubleComplex(0,0);
for(b_mm=0;b_mm<19;b_mm++)
{
if(b_mm<2*ll+1)
{
xt2 = make_cuDoubleComplex(0,0);
for(ju=0;ju<natoms;ju++)
{
if(h != ju)
{
dx[(h*147)+ju]=(x11[h]-x11[ju]);
dy[(h*147)+ju]=(y11[h]-y11[ju]);
dz[(h*147)+ju]=(z11[h]-z11[ju]);
r[(h*147)+ju]=sqrtf(dx[(h*147)+ju]*dx[(h*147)+ju] + dy[(h*147)+ju]*dy[(h*147)+ju] + dz[(h*147)+ju]*dz[(h*147)+ju]);
fophi1[(h*147)+ju]= atan2f(dy[(h*147)+ju],dx[(h*147)+ju]);
if(r[(h*147)+ju] <= rcut)
{
rcutfun[(h*147)+ju] = 0.5*(cosf(pi*r[(h*147)+ju]*rcut_inv)+1);
}
else
{
rcutfun[(h*147)+ju] = 0.0;
}
func1[(h*147)+ju] = expf(-netaa[n]*r[(h*147)+ju]*r[(h*147)+ju]);
term1[(h*147)+ju] = func1[(h*147)+ju]*rcutfun[(h*147)+ju];
G_p[(h*147)+ju] = term1[(h*147)+ju];
mm = -ll+b_mm;
m_1 = abs(mm);
if (mm<0)
{
temp1 = cuCmul(q, make_cuDoubleComplex(-m_1,0));
ex1[(h*147)+ju] = cuCmul(temp1, make_cuDoubleComplex(fophi1[(h*147)+ju],0));
comp_1_w[(h*147)+ju] = my_complex_exp(ex1[(h*147)+ju]);
}
else
{
temp2 = cuCmul(q, make_cuDoubleComplex(m_1,0));
ex2[(h*147)+ju] = cuCmul(temp2, make_cuDoubleComplex(fophi1[(h*147)+ju],0));
ex3[(h*147)+ju] = my_complex_exp (ex2[(h*147)+ju]);
float result = 1.0;
int zz;
for(zz=9; zz>=1; zz--)
{
if(zz<=mm)
{
result *= (-1.0);
}
}
comp_1_w[(h*147)+ju] = cuCmul(ex3[(h*147)+ju], make_cuDoubleComplex(result, 0));
//printf("comp_1_w values : %f %f\n",comp_1_w[(h*147)+ju].x,comp_1_w[(h*147)+ju].y);
}
//mult1[(h*147)+ju] = cuCmul(comp_1_w[(h*147)+ju], make_cuDoubleComplex(cons1[(ll*10)+b_mm], 0));
mult1[(h*147)+ju] = comp_1_w[(h*147)+ju];
//printf("b_mm,ll,cons1 values : %d\t%d\t%f\n",b_mm,ll,cons1[1]);
//printf("mult1 values : %f %f\n",mult1[(h*147)+ju].x,mult1[(h*147)+ju].y);
Lmmnn = Lmn_all[ju][ll][m_1];
//printf("Lmn_all : %f\n",Lmn_all[ju][ll][m_1]);
//printf("Lmmnn : %f\n",Lmmnn);
useY1[(h*147)+ju] = cuCmul(mult1[(h*147)+ju], make_cuDoubleComplex(Lmmnn, 0));
//printf("useY1 values : %f %f\n",useY1[(h*147)+ju].x,useY1[(h*147)+ju].y);
mult4[(h*147)+ju] = cuCmul(useY1[(h*147)+ju], make_cuDoubleComplex(G_p[(h*147)+ju], 0));
//printf("mult4 values : %f %f\n",mult4[(h*147)+ju].x,mult4[(h*147)+ju].y);
xt2 = cuCadd(xt2,mult4[(h*147)+ju]);
//xt2 = cuCadd(make_cuDoubleComplex(1,0),mult4[(h*147)+ju]);
//printf("xt2 values : %f %f\n",xt2.x,xt2.y);
}
}
xt6[(h*147)+ju] = cuConj(xt2);
temp3[(h*147)+ju] = cuCmul(xt2,xt6[(h*147)+ju]);
p_ch_p[(h*147)+ju] = cuCadd(p_ch_p[(h*147)+ju],temp3[(h*147)+ju]);
//printf("p_ch_p values : %f %f\n",p_ch_p[(h*147)+ju].x,p_ch_p[(h*147)+ju].y);
tot_cnlm_sum_p[h][n][ll][b_mm] = xt2;
printf("tot_cnlm_sum_p values : %f %f\n",tot_cnlm_sum_p[h][n][ll][b_mm].x,tot_cnlm_sum_p[h][n][ll][b_mm].y);
}
}
}
}
}
//printf("Lmmnn value : %f\n",Lmmnn);
}
int main()
{
float *dx = (float*)malloc((147*147)*sizeof(float));
float *dy = (float*)malloc((147*147)*sizeof(float));
float *dz = (float*)malloc((147*147)*sizeof(float));
float *x11 = (float*)malloc(147*sizeof(float));
float *y11 = (float*)malloc(147*sizeof(float));
float *z11 = (float*)malloc(147*sizeof(float));
float *r = (float*)malloc((147*147)*sizeof(float));
float *pcolhea1 = (float*)malloc((147*147)*sizeof(float));
float *fophi1 = (float*)malloc((147*147)*sizeof(float));
float *rcutfun = (float*)malloc((147*147)*sizeof(float));
float *term1 = (float*)malloc((147*147)*sizeof(float));
float *G_p = (float*)malloc((147*147)*sizeof(float));
float *func1 = (float*)malloc((147*147)*sizeof(float));
cuDoubleComplex *ex1 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *comp_1_w = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *ex2 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *ex3 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *mult1 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *useY1 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *mult4 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
//cuDoubleComplex *xt2 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *xt6 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *temp3 = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
cuDoubleComplex *p_ch_p = (cuDoubleComplex*)malloc((147*147)*sizeof(cuDoubleComplex));
float *cons11= (float*)malloc(100*sizeof(float));
float *cons1 = (float*)malloc(210*sizeof(float));
FILE * fp1;
fp1 = fopen ("cons11.txt", "r");
int index = 0;
while(fscanf(fp1, "%f",(cons11+index))==1)
index++;
fclose(fp1);
int l,m;
for (l=0;l<10;l++)
{
for(m=0;m<21;m++)
{
if(m<2*l+1)
{
cons1[(l*10)+m] = cons11[(l*l)+m];
}
else
{
cons1[(l*10)+m] = 0.0;
}
}
}
float *devdx, *d_x11, *devdy, *d_y11, *devdz, *d_z11, *d_r, *d_pcolhea1, *d_fophi1, *d_rcutfun, *d_term1, *d_G_p, *d_func1,*d_cons1;
cuDoubleComplex *d_ex1, *d_ex2, *d_ex3, *d_comp_1_w, *d_mult1, *d_useY1, *d_mult4, *d_xt6, *d_temp3, *d_p_ch_p;
cudaMalloc(&devdx,(147*147)*sizeof(float));
cudaMalloc(&d_x11,(147)*sizeof(float));
cudaMalloc(&devdy,(147*147)*sizeof(float));
cudaMalloc(&d_y11,(147)*sizeof(float));
cudaMalloc(&devdz,(147*147)*sizeof(float));
cudaMalloc(&d_z11,(147)*sizeof(float));
cudaMalloc(&d_r,(147*147)*sizeof(float));
cudaMalloc(&d_pcolhea1,(147*147)*sizeof(float));
cudaMalloc(&d_fophi1,(147*147)*sizeof(float));
cudaMalloc(&d_rcutfun,(147*147)*sizeof(float));
cudaMalloc(&d_term1,(147*147)*sizeof(float));
cudaMalloc(&d_G_p,(147*147)*sizeof(float));
cudaMalloc(&d_func1,(147*147)*sizeof(float));
cudaMalloc(&d_ex1,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_ex2,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_ex3,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_comp_1_w,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_mult1,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_useY1,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_cons1,(210)*sizeof(float));
cudaMalloc(&d_mult4,(147*147)*sizeof(cuDoubleComplex));
//cudaMalloc(&d_xt2,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_xt6,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_temp3,(147*147)*sizeof(cuDoubleComplex));
cudaMalloc(&d_p_ch_p,(147*147)*sizeof(cuDoubleComplex));
//float x11[4];
int h;
FILE *fp;
if ((fp = fopen("XYZ1.txt","r")) == NULL){
printf("Error! opening file");
exit(1);
}
for(h=0;h<147;h++){
fscanf(fp, "%f", (x11+h));
}
for(h=147;h<294;h++)
{
fscanf(fp, "%f", (y11+h-147));
}
for(h=294;h<441;h++)
{
fscanf(fp, "%f", (z11+h-294));
}
cudaMemcpy(d_x11, x11, (147*sizeof(float)), cudaMemcpyHostToDevice);
cudaMemcpy(d_y11, y11, (147*sizeof(float)), cudaMemcpyHostToDevice);
cudaMemcpy(d_z11, z11, (147*sizeof(float)), cudaMemcpyHostToDevice);
cudaMemcpy(d_cons1, cons1, (210*sizeof(float)), cudaMemcpyHostToDevice);
//cudaMemcpy(d_r, r, (4*sizeof(float)), cudaMemcpyHostToDevice);
dim3 blocksize(147,147);
dim3 threadsize(1,1);
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
fc<<<blocksize,threadsize>>>(d_x11,devdx,d_y11,devdy,d_z11,devdz,d_r,d_pcolhea1,d_fophi1,d_rcutfun,d_term1,d_G_p,d_func1,d_ex1,d_comp_1_w,d_ex2,d_ex3,d_mult1,d_cons1,d_useY1,d_mult4,d_xt6,d_temp3,d_p_ch_p);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsed, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
//printf("The elapsed time in gpu was %.2f ms\n", elapsed);
//cudaDeviceSynchronize();
checkCudaErr(cudaDeviceSynchronize(), "Synchronization");
checkCudaErr(cudaGetLastError(), "GPU");
checkCudaErr(cudaPeekAtLastError(), "GPUpeek");
cudaError err1 = cudaMemcpy(dx, devdx,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err2 = cudaMemcpy(dy, devdy,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err3 = cudaMemcpy(dz, devdz,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err4 = cudaMemcpy(r, d_r,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err5 = cudaMemcpy(pcolhea1, d_pcolhea1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err6 = cudaMemcpy(fophi1, d_fophi1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err7 = cudaMemcpy(rcutfun, d_rcutfun,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err8 = cudaMemcpy(term1, d_term1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err9 = cudaMemcpy(G_p, d_G_p,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err10 =cudaMemcpy(func1, d_func1,((147*147)*sizeof(float)),cudaMemcpyDeviceToHost);
cudaError err11 =cudaMemcpy(ex1, d_ex1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaError err12 =cudaMemcpy(comp_1_w, d_comp_1_w,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaError err13 = cudaMemcpy(mult1, d_mult1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaError err14 = cudaMemcpy(useY1, d_useY1,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(mult4, d_mult4,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(ex2, d_ex2,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(ex3, d_ex3,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
//cudaMemcpy(xt2, d_xt2,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(xt6, d_xt6,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(temp3, d_temp3,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
cudaMemcpy(p_ch_p, d_p_ch_p,((147*147)*sizeof(cuDoubleComplex)),cudaMemcpyDeviceToHost);
if((err1 || err2 || err3 || err4 || err5 || err6 || err7 || err8 || err9 || err10 || err11 || err12 || err13 || err14)!=cudaSuccess){
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err1));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err2));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err3));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err4));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err5));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err6));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err7));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err8));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err9));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err10));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err11));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err12));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err13));
printf("CUDA error copying to Host : %s\n",cudaGetErrorString(err14));
}
cudaEventSynchronize(end);
return 0;
}
Could you please guide what may be the possible reason for this error. Also when I do these calculations for 2 points rather than 147 points, I got the results for 2 points. But ultimately I have to do it for 147 points.
Thank You and have a nice day!
- If you don’t wish to employ and demonstrate rigorous, proper CUDA error checking, I won’t be able to spend any time on your questions.
- If you ask for help with codes that require input files to run, I probably won’t be able to help with those either.