Add more apps to 2_Cookbook

Change-Id: Iafe462df9726a32f450bd240a2de3eaa73a10057

samples/2_Cookbook/3_shared_memory/sharedMemory.cpp

@@ -26,13 +26,12 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 
 
-#define WIDTH     1024
-#define HEIGHT    1024
+#define WIDTH     64
 
-#define NUM       (WIDTH*HEIGHT)
+#define NUM       (WIDTH*WIDTH)
 
-#define THREADS_PER_BLOCK_X  16
-#define THREADS_PER_BLOCK_Y  16
+#define THREADS_PER_BLOCK_X  4
+#define THREADS_PER_BLOCK_Y  4
 #define THREADS_PER_BLOCK_Z  1
 
 // Device (Kernel) function, it must be void
@@ -40,15 +39,14 @@ THE SOFTWARE.
 __global__ void matrixTranspose(hipLaunchParm lp,
                                 float *out,
                                 float *in,
-                                const int width,
-                                const int height)
+                                const int width)
 {
-    __shared__ float sharedMem[16*16];
+    __shared__ float sharedMem[WIDTH*WIDTH];
 
     int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
     int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
 
-    sharedMem[y * width + x] = in[x * height + y];
+    sharedMem[y * width + x] = in[x * width + y];
 
     __syncthreads();
 
@@ -59,14 +57,13 @@ __global__ void matrixTranspose(hipLaunchParm lp,
 void matrixTransposeCPUReference(
     float * output,
     float * input,
-    const unsigned int width,
-    const unsigned int height)
+    const unsigned int width)
 {
-    for(unsigned int j=0; j < height; j++)
+    for(unsigned int j=0; j < width; j++)
     {
         for(unsigned int i=0; i < width; i++)
         {
-            output[i*height + j] = input[j*width + i];
+            output[i*width + j] = input[j*width + i];
         }
     }
 }
@@ -106,22 +103,22 @@ int main() {
 
   // Lauching kernel from host
   hipLaunchKernel(matrixTranspose,
-                  dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
+                  dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                   0, 0,
-                  gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
+                  gpuTransposeMatrix , gpuMatrix, WIDTH);
 
   // Memory transfer from device to host
   hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
 
   // CPU MatrixTranspose computation
-  matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH, HEIGHT);
+  matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
 
   // verify the results
   errors = 0;
   double eps = 1.0E-6;
   for (i = 0; i < NUM; i++) {
-    if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
+    if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
     printf("%d cpu: %f gpu  %f\n",i,cpuTransposeMatrix[i],TransposeMatrix[i]);
       errors++;
     }