Add more apps to 2_Cookbook

Change-Id: Iafe462df9726a32f450bd240a2de3eaa73a10057


[ROCm/hip-tests commit: 04af19866f]

projects/hip-tests/samples/2_Cookbook/3_shared_memory/Makefile

@@ -10,7 +10,7 @@ TARGET=hcc
 SOURCES = sharedMemory.cpp
 OBJECTS = $(SOURCES:.cpp=.o)
 
-EXECUTABLE=./exe
+EXECUTABLE=./sharedMemory
 
 .PHONY: test
 

projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md

@@ -0,0 +1,42 @@
+## Using shared memory ###
+
+Earlier we learned how to write our first hip program, in which we compute Matrix Transpose. In this tutorial, we'll explain how to use the shared memory to improve the performance.
+
+## Introduction:
+
+As we mentioned earlier  that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use static shared memory and will explain the dynamic one latter.
+
+## Requirement:
+For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) 
+
+## prerequiste knowledge:
+
+Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
+
+## Simple Matrix Transpose 
+
+We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
+
+## Shared Memory
+
+Shared memory is way more faster than that of global and constant memory and accessible to all the threads in the block. If the size of shared memory is known at compile time, we can specify the size and will use the static shared memory. In the same sourcecode, we will use the `__shared__` variable type qualifier as follows:
+
+`  __shared__ float sharedMem[1024*1024];`
+
+Be careful while using shared memory, since all threads within the block can access the shared memory, we need to sync the operation of individual threads by using:
+
+`  __syncthreads();`
+
+## How to build and run:
+Use the make command and execute it using ./exe
+Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.
+
+## More Info:
+- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md)
+- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md)
+- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
+- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
+- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
+- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
+- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
+- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)

projects/hip-tests/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++;
     }