add ch4

[book_gpu.git] / BookGPU / Chapters / chapter2 / ex3.cu

diff --git a/BookGPU/Chapters/chapter2/ex3.cu b/BookGPU/Chapters/chapter2/ex3.cu
index fbdf3a24d3d0b8f70975712be40383111e4bcab0..cddcc309d41329f5e440005c31a6f4bf9ddb677a 100644 (file)
--- a/BookGPU/Chapters/chapter2/ex3.cu
+++ b/BookGPU/Chapters/chapter2/ex3.cu
@@ -6,16 +6,12 @@
 #include "cutil_inline.h"
 #include <cublas_v2.h>
 
 #include "cutil_inline.h"
 #include <cublas_v2.h>
 

-

 const int width=16;
 const int nbTh=width*width;
 
 const int size=1024;
 const  int sizeMat=size*size;
 
 const int width=16;
 const int nbTh=width*width;
 
 const int size=1024;
 const  int sizeMat=size*size;
 

-
-
-

 __global__ 
 void matmul(float *d_A, float *d_B, float *d_C) {
        int i= blockIdx.y*blockDim.y+ threadIdx.y;
 __global__ 
 void matmul(float *d_A, float *d_B, float *d_C) {
        int i= blockIdx.y*blockDim.y+ threadIdx.y;


@@ -26,15 +22,10 @@ void matmul(float *d_A, float *d_B, float *d_C) {
                sum+=d_A[i*size+k]*d_B[k*size+j];
        }       
        d_C[i*size+j]=sum;
                sum+=d_A[i*size+k]*d_B[k*size+j];
        }       
        d_C[i*size+j]=sum;

-

 }
 
 }
 

-
-
-

 int main( int argc, char** argv) 
 {
 int main( int argc, char** argv) 
 {

-

        float *h_arrayA=(float*)malloc(sizeMat*sizeof(float));
        float *h_arrayB=(float*)malloc(sizeMat*sizeof(float));
        float *h_arrayC=(float*)malloc(sizeMat*sizeof(float));
        float *h_arrayA=(float*)malloc(sizeMat*sizeof(float));
        float *h_arrayB=(float*)malloc(sizeMat*sizeof(float));
        float *h_arrayC=(float*)malloc(sizeMat*sizeof(float));


@@ -46,9 +37,7 @@ int main( int argc, char** argv)
        cudaMalloc((void**)&d_arrayB,sizeMat*sizeof(float));
        cudaMalloc((void**)&d_arrayC,sizeMat*sizeof(float));
 
        cudaMalloc((void**)&d_arrayB,sizeMat*sizeof(float));
        cudaMalloc((void**)&d_arrayC,sizeMat*sizeof(float));
 

-

        srand48(32);
        srand48(32);

-

        for(int i=0;i<sizeMat;i++) {
                h_arrayA[i]=drand48();
                h_arrayB[i]=drand48();
        for(int i=0;i<sizeMat;i++) {
                h_arrayA[i]=drand48();
                h_arrayB[i]=drand48();


@@ -61,7 +50,6 @@ int main( int argc, char** argv)
        cudaMemcpy(d_arrayB,h_arrayB, sizeMat * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(d_arrayC,h_arrayC, sizeMat * sizeof(float), cudaMemcpyHostToDevice);
 
        cudaMemcpy(d_arrayB,h_arrayB, sizeMat * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(d_arrayC,h_arrayC, sizeMat * sizeof(float), cudaMemcpyHostToDevice);
 

-

        unsigned int timer_cpu = 0;
        cutilCheckError(cutCreateTimer(&timer_cpu));
   cutilCheckError(cutStartTimer(timer_cpu));
        unsigned int timer_cpu = 0;
        cutilCheckError(cutCreateTimer(&timer_cpu));
   cutilCheckError(cutStartTimer(timer_cpu));


@@ -77,20 +65,13 @@ int main( int argc, char** argv)
        printf("CPU processing time : %f (ms) \n", cutGetTimerValue(timer_cpu));
        cutDeleteTimer(timer_cpu);
 
        printf("CPU processing time : %f (ms) \n", cutGetTimerValue(timer_cpu));
        cutDeleteTimer(timer_cpu);
 

-
-
-

        unsigned int timer_gpu = 0;
        cutilCheckError(cutCreateTimer(&timer_gpu));
   cutilCheckError(cutStartTimer(timer_gpu));
 
        unsigned int timer_gpu = 0;
        cutilCheckError(cutCreateTimer(&timer_gpu));
   cutilCheckError(cutStartTimer(timer_gpu));
 

-
-

        dim3 dimGrid(size/width,size/width);
        dim3 dimBlock(width,width);
 
        dim3 dimGrid(size/width,size/width);
        dim3 dimBlock(width,width);
 

-       printf("%d %d\n",dimGrid.x,dimBlock.x);
-

        matmul<<<dimGrid,dimBlock>>>(d_arrayA,d_arrayB,d_arrayC);
        cudaThreadSynchronize();
        
        matmul<<<dimGrid,dimBlock>>>(d_arrayA,d_arrayB,d_arrayC);
        cudaThreadSynchronize();
        


@@ -100,12 +81,10 @@ int main( int argc, char** argv)
        
        cudaMemcpy(h_arrayCgpu,d_arrayC, sizeMat * sizeof(float), cudaMemcpyDeviceToHost);
        
        
        cudaMemcpy(h_arrayCgpu,d_arrayC, sizeMat * sizeof(float), cudaMemcpyDeviceToHost);
        

-       int good=1;

        for(int i=0;i<sizeMat;i++)
                if (fabs(h_arrayC[i]-h_arrayCgpu[i])>1e-4)
                        printf("%f %f\n",h_arrayC[i],h_arrayCgpu[i]);
        
        for(int i=0;i<sizeMat;i++)
                if (fabs(h_arrayC[i]-h_arrayCgpu[i])>1e-4)
                        printf("%f %f\n",h_arrayC[i],h_arrayCgpu[i]);
        

-

        cudaFree(d_arrayA);
        cudaFree(d_arrayB);
        cudaFree(d_arrayC);
        cudaFree(d_arrayA);
        cudaFree(d_arrayB);
        cudaFree(d_arrayC);


@@ -113,7 +92,5 @@ int main( int argc, char** argv)
        free(h_arrayB);
        free(h_arrayC);
        free(h_arrayCgpu);
        free(h_arrayB);
        free(h_arrayC);
        free(h_arrayCgpu);

-

        return 0;
        return 0;

-

 }
 }