测试求和 waiting

//duide 1

#include "h.h"
#define  G 4 //灰度阶
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
//测试求和
float ComputeFeature1( int* pdMatrix,int total)
{
 int i,j;
 float f1=0;
 for(i=0; i<G; i++)
 {
  for(j=0; j<G; j++)
  {
   float t;
   t=(float)pdMatrix[i*G+j]/(float)total;
   f1+= (t)*(t);
  }
 }
 return f1;
}
__global__  static void ComputeFeature1( float *sum1,int* pdMatrix,int width,int height,int *total)  //186ye
{
 int i;
 float t;
 int tid=threadIdx.x+threadIdx.y*4;//
 int tid_in_x = threadIdx.x;
 int tid_in_y = threadIdx.y;
 __shared__ float stemp1[G*G];

 if(tid_in_y<G &&(tid_in_x<G))
 {
 stemp1[tid_in_x*G+tid_in_y]=0;
 }

 __syncthreads();
 //归一化
 if(tid_in_y<G &&(tid_in_x<G))
 {
  stemp1[tid_in_x*G+tid_in_y]=(float)pdMatrix[tid]/((float)(*total));
 }
  __syncthreads();
if(threadIdx.x<G && threadIdx.y<G)
 {
  stemp1[threadIdx.x*G+threadIdx.y]=stemp1[threadIdx.x*G+threadIdx.y]*stemp1[threadIdx.x*G+threadIdx.y];
  
    }
 __syncthreads();
 
if(tid<G*G)
{
 for(i=G*G/2; i>0; i/=2)
 {
        if(tid<i)
  {
            stemp1[tid] = stemp1[tid] + stemp1[tid+i] ;

  }
        __syncthreads();

    }
}
  __syncthreads();
 *sum1 = stemp1[0];
DUI
//if(tid<G*G)
//{
// for(i=G*G/2; i>0; i/=2)
// {
//        if(tid<i)
//  {
//            pdMatrix[tid] = pdMatrix[tid] + pdMatrix[tid+i] ;
//
//  }
//        __syncthreads();
//
//    }
//}
//  __syncthreads();
// *sum1 = pdMatrix[0];
//

}
int main()
{
 float *sum1,sum2;
 //dim3 grid(2,2);
 dim3 block(4,4);
 int* pdMatrix;
 int*p;
 float *g_f1,*g_f2;
 int width,height;
 width=height=4;
 int t=height*width;
 //int tt=height*width/4;
 p=(int *)malloc(height*width* sizeof(float));
 int tot=0;
 int *total;
 float f1=0;
 cudaMalloc((void**)&total, sizeof(int));
 for(int i=0;i<t;i++)
 {
  p[i]=i;//i;
  tot+=p[i];
 }
 f1=ComputeFeature1( p,tot);
 printf("%f\n",f1);
 //tot=1;
 cudaMalloc((void**)&total, sizeof(int));
 cudaMemset(total,0,sizeof(int));
 
 cudaMemcpy(total,&tot,sizeof(int), cudaMemcpyHostToDevice);

 cudaMalloc((void**)&pdMatrix,height*width* sizeof(int));
 //cudaMalloc((void**)&g_f1,tt* sizeof(float));
 //cudaMemset(g_f1,0,sizeof(float)*tt);
 cudaMalloc((void**)&sum1, sizeof(float));
 cudaMemset(sum1,0,sizeof(float));
 //g_f2=(float*)malloc(tt* sizeof(float));
 //memset(g_f2,0,sizeof(float)*tt);
 sum2=0;
    cudaMemcpy(pdMatrix,p,height*width*sizeof(int), cudaMemcpyHostToDevice);
 ComputeFeature1<<<1,block>>>(sum1,pdMatrix,width,height,total);
 cudaError_t cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

  //cudaMemcpy(g_f2,g_f1,tt*sizeof(float), cudaMemcpyDeviceToHost);
   cudaMemcpy(&sum2,sum1,sizeof(float), cudaMemcpyDeviceToHost);
  cudaStatus = cudaDeviceSynchronize();
 if (cudaStatus != cudaSuccess) {
  fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);

 }
 printf("%f\n",sum2);
 //printf("\n%d\n",sum2);
 
 Error:
   
    return cudaStatus;
}

 

 

//duide 1 end

 

#include "h.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
//测试求和

__global__  static void ComputeFeature1( int *sum1,int* pdMatrix,int width,int height,float *g_f1)  //186ye
{
 //pdMatrix[threadIdx.x]=pdMatrix[threadIdx.x]*pdMatrix[threadIdx.x];

 __shared__ float stemp1[G][G];
 int i,j;
 unsigned int xIndex=__mul24(blockDim.x,blockIdx.x)+threadIdx.x;
 unsigned int yIndex=__mul24(blockDim.y,blockIdx.y)+threadIdx.y;

  unsigned int index=yIndex*width+xIndex; //?
 int tid_in_x=threadIdx.x;
 int tid_in_y=threadIdx.y;
 if(tid_in_x<G &&tid_in_y<G){
  stemp1[tid_in_x][tid_in_y]=0.0f;}
 __syncthreads();
 if((index<width*height)&&(threadIdx.x<G)&&(threadIdx.y<G)) //index_in
 {
  stemp1[threadIdx.x][threadIdx.y]=pdMatrix[index];
 } 
 __syncthreads();

 //float sum1;
 //sum1=0.0f;
 //for(i=0;i<16;i++)
 //{
 // //for(j=0;j<G;j++)
 // //{
 //          *sum1+= pdMatrix[i];
 // //}
 //}

//   这种对
if(index<height*width)
{
 for(i=height*width/2; i>0; i/=2)
 {
        if(index<i)
  {
            pdMatrix[index] = pdMatrix[index] + pdMatrix[index+i] ;

  }
        __syncthreads();

    }
}
  __syncthreads();
 *sum1 = pdMatrix[0];

 ///

/if(sum1!=0)
 //{
 //if(threadIdx.x<G &&threadIdx.y<G){
 stemp1[threadIdx.x][threadIdx.y]=stemp1[threadIdx.x][threadIdx.y]/sum1;
 //}
 //__syncthreads();
 //if(threadIdx.x<G &&threadIdx.y<G){
 stemp1[threadIdx.x][threadIdx.y]=stemp1[threadIdx.x][threadIdx.y]*stemp1[threadIdx.x][threadIdx.y];
 //}
 //__syncthreads();

//这种不对？？？？？？？？？？？？？？
 /*if(threadIdx.x<G &&threadIdx.y<G){
 for(j=blockDim.y*2;j>0;j--)
 {
  if(threadIdx.y<j)
  {
   stemp1[threadIdx.x][threadIdx.y]+=stemp1[threadIdx.x][threadIdx.y+j];
  }
  __syncthreads();
 }
 for( i=blockDim.x*2;i > 0 ;i--)
 {
  if (threadIdx.y < j )
  {
   stemp1[threadIdx.x][0]+= stemp1[threadIdx.x+i][0];
  }
  __syncthreads();
 }
 }
 if( threadIdx.x==0&&threadIdx.y==0){
  g_f1[blockIdx.x*2+blockIdx.y]=stemp1[0][0];//

 //}
}*/
}
int main()
{
 int *sum1,sum2;
 dim3 grid(2,2,1);
 dim3 block(2,2,1);
 int* pdMatrix;
 int*p;
 float *g_f1,*g_f2;
 int width,height;
 width=height=4;
 int t=height*width;
 int tt=height*width/4;
 p=(int *)malloc(height*width* sizeof(float));
 for(int i=0;i<t;i++)
 {
  p[i]=i;//i;
 }
 
 cudaMalloc((void**)&pdMatrix,height*width* sizeof(int));
 cudaMalloc((void**)&g_f1,tt* sizeof(float));
 cudaMemset(g_f1,0,sizeof(float)*tt);
 cudaMalloc((void**)&sum1, sizeof(int));
 cudaMemset(sum1,0,sizeof(int));
 g_f2=(float*)malloc(tt* sizeof(float));
 memset(g_f2,0,sizeof(float)*tt);
 sum2=0;
    cudaMemcpy(pdMatrix,p,height*width*sizeof(int), cudaMemcpyHostToDevice);
 ComputeFeature1<<<grid,block>>>(sum1,pdMatrix,width,height,g_f1);
 cudaError_t cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

  cudaMemcpy(g_f2,g_f1,tt*sizeof(float), cudaMemcpyDeviceToHost);
   cudaMemcpy(&sum2,sum1,sizeof(int), cudaMemcpyDeviceToHost);
  cudaStatus = cudaDeviceSynchronize();
 if (cudaStatus != cudaSuccess) {
  fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);

 }
 printf("%d\n",sum2);sum2=0;
 for(int i=0;i<tt;i++)
 {
  printf("%f ",g_f2[i]);
   sum2+=g_f2[i];
 }
 printf("\n%d\n",sum2);

 //p[width*height]={0,0,0,1,0,0,1,0,1,1,0,0,0,1,0,1};
 
 Error:
   
    return cudaStatus;
}