Init changed to reduce magnitude of the L2 norm of the kernel results. Updated sample_output

Ram3sh · Ram3sh · commit 3c024623aa18 · 2021-03-11T16:28:50.000-08:00
diff --git a/sample_output/OUT_WHAMO_CUDA b/sample_output/OUT_WHAMO_CUDA
@@ -2,14 +2,14 @@ srun -n 1 -p nvidia bin/sw4ck ../src/sw4ck.in
 Reading from file ../src/sw4ck.in
 Launching sw4 kernels
 
-Kernel 1 time 4.69094
-Kernel 2 time 1.25952
-Kernel 3 time 1.26464
-Kernel 4 time 1.11923
-Kernel 5 time 5.38419
+Kernel 1 time 4.72678
+Kernel 2 time 1.26259
+Kernel 3 time 1.26054
+Kernel 4 time 1.11821
+Kernel 5 time 5.41491
 
-Total kernel runtime = 14
+Total kernel runtime = 13
 
-Norm of output 0x1.b5fa52fe0079dp+59
-Norm of output 986238393426103936
-Error = -1.3e-14 %
+Norm of output 0x1.941a40aec142ep+7
+Norm of output 202.0512747393526638
+Error = 0 %
diff --git a/sample_output/OUT_WHAMO_CUDA_RAJA b/sample_output/OUT_WHAMO_CUDA_RAJA
@@ -2,14 +2,14 @@ srun -n 1 -p nvidia bin/sw4ck ../src/sw4ck.in
 Reading from file ../src/sw4ck.in
 Launching sw4 kernels
 
-Kernel 1 time 6.92634
-Kernel 2 time 1.70496
-Kernel 3 time 1.74182
-Kernel 4 time 1.36806
-Kernel 5 time 8.98355
+Kernel 1 time 6.94067
+Kernel 2 time 1.7152
+Kernel 3 time 1.69062
+Kernel 4 time 1.33939
+Kernel 5 time 8.64973
 
-Total kernel runtime = 21
+Total kernel runtime = 20
 
-Norm of output 0x1.b5fa52fe0079fp+59
-Norm of output 986238393426104192
-Error = 1.3e-14 %
+Norm of output 0x1.941a40aec142ep+7
+Norm of output 202.0512747393526638
+Error = 0 %
diff --git a/sample_output/OUT_WHAMO_HIP_CC b/sample_output/OUT_WHAMO_HIP_CC
@@ -0,0 +1,15 @@
+srun -n 1 -p amd ./sw4ck sw4ck.in
+Reading from file sw4ck.in
+Launching sw4 kernels
+
+Kernel 1 time 58.0426
+Kernel 2 time 24.1525
+Kernel 3 time 23.6964
+Kernel 4 time 17.5566
+Kernel 5 time 50.5733
+
+Total kernel runtime = 174
+
+Norm of output 0x1.941a40ae9d6fap+7
+Norm of output 202.05127473518206216
+Error = -2.1e-09 %
diff --git a/sample_output/OUT_WHAMO_HIP_HIPCC b/sample_output/OUT_WHAMO_HIP_HIPCC
@@ -0,0 +1,15 @@
+srun -n 1 -p amd ./sw4ck sw4ck.in
+Reading from file sw4ck.in
+Launching sw4 kernels
+
+Kernel 1 time 15.3942
+Kernel 2 time 4.6632
+Kernel 3 time 3.43617
+Kernel 4 time 3.74249
+Kernel 5 time 15.4687
+
+Total kernel runtime = 42
+
+Norm of output 0x1.941a40aec142ep+7
+Norm of output 202.0512747393526638
+Error = 0 %
diff --git a/sample_output/OUT_WHAMO_HIP_RAJA b/sample_output/OUT_WHAMO_HIP_RAJA
@@ -2,14 +2,14 @@ srun -n 1 -p amd bin/sw4ck ../src/sw4ck.in
 Reading from file ../src/sw4ck.in
 Launching sw4 kernels
 
-Kernel 1 time 21.2905
-Kernel 2 time 30.6422
-Kernel 3 time 37.4966
-Kernel 4 time 28.6283
-Kernel 5 time 22.5596
+Kernel 1 time 21.0319
+Kernel 2 time 30.7601
+Kernel 3 time 37.5654
+Kernel 4 time 28.6686
+Kernel 5 time 22.607
 
 Total kernel runtime = 140
 
-Norm of output 0x1.b5fa52fe0079fp+59
-Norm of output 986238393426104192
-Error = 1.3e-14 %
+Norm of output 0x1.941a40aec142ep+7
+Norm of output 202.0512747393526638
+Error = 0 %
diff --git a/sample_output/init/OUT_CORONA b/sample_output/init/OUT_CORONA
diff --git a/sample_output/init/OUT_CORONA_RAJA b/sample_output/init/OUT_CORONA_RAJA
diff --git a/sample_output/init/OUT_LASSEN b/sample_output/init/OUT_LASSEN
diff --git a/sample_output/init/OUT_LASSEN_RAJA b/sample_output/init/OUT_LASSEN_RAJA
diff --git a/sample_output/init/OUT_WHAMO_CUDA b/sample_output/init/OUT_WHAMO_CUDA
@@ -0,0 +1,15 @@
+srun -n 1 -p nvidia bin/sw4ck ../src/sw4ck.in
+Reading from file ../src/sw4ck.in
+Launching sw4 kernels
+
+Kernel 1 time 4.69094
+Kernel 2 time 1.25952
+Kernel 3 time 1.26464
+Kernel 4 time 1.11923
+Kernel 5 time 5.38419
+
+Total kernel runtime = 14
+
+Norm of output 0x1.b5fa52fe0079dp+59
+Norm of output 986238393426103936
+Error = -1.3e-14 %
diff --git a/sample_output/init/OUT_WHAMO_CUDA_RAJA b/sample_output/init/OUT_WHAMO_CUDA_RAJA
@@ -0,0 +1,15 @@
+srun -n 1 -p nvidia bin/sw4ck ../src/sw4ck.in
+Reading from file ../src/sw4ck.in
+Launching sw4 kernels
+
+Kernel 1 time 6.92634
+Kernel 2 time 1.70496
+Kernel 3 time 1.74182
+Kernel 4 time 1.36806
+Kernel 5 time 8.98355
+
+Total kernel runtime = 21
+
+Norm of output 0x1.b5fa52fe0079fp+59
+Norm of output 986238393426104192
+Error = 1.3e-14 %
diff --git a/sample_output/init/OUT_WHAMO_HIP b/sample_output/init/OUT_WHAMO_HIP
diff --git a/sample_output/init/OUT_WHAMO_HIP_RAJA b/sample_output/init/OUT_WHAMO_HIP_RAJA
@@ -0,0 +1,15 @@
+srun -n 1 -p amd bin/sw4ck ../src/sw4ck.in
+Reading from file ../src/sw4ck.in
+Launching sw4 kernels
+
+Kernel 1 time 21.2905
+Kernel 2 time 30.6422
+Kernel 3 time 37.4966
+Kernel 4 time 28.6283
+Kernel 5 time 22.5596
+
+Total kernel runtime = 140
+
+Norm of output 0x1.b5fa52fe0079fp+59
+Norm of output 986238393426104192
+Error = 1.3e-14 %
diff --git a/src/configs/make.CC b/src/configs/make.CC
@@ -12,11 +12,11 @@ HSA_ROOT_DIR = /opt/rocm-4.0.1/hsa
 
 # --save-temps aborts compilation. -fno-cray fixes this. Incorrect and slow results with CC. Correct slower results wth CC -fno-cray
 # Add -fsave-loopmark for optimization report
-MORE_FLAGS = -I${ROCM_PATH}/include -O3  -x hip -Winconsistent-missing-override  -DNO_RAJA=1 -Wall -Wno-unknown-pragmas --rocm-path=/opt/rocm-4.0.1 --offload-arch=gfx906 -fsave-loopmark -fno-cray
+MORE_FLAGS = -I${ROCM_PATH}/include -Ofast  -x hip -Winconsistent-missing-override  -Wall -Wno-unknown-pragmas --rocm-path=/opt/rocm-4.0.1 --offload-arch=gfx906 -fsave-loopmark -fno-cray 
 
 LINKFLAGS =  -std=c++11 
 
-EXTRA_CXX_FLAGS =  -I$(RAJA_LOCATION)/include  -DENABLE_HIP $(MORE_FLAGS) -Winconsistent-missing-override 
+EXTRA_CXX_FLAGS =  -I$(RAJA_LOCATION)/include  -DENABLE_HIP $(MORE_FLAGS) -DNO_RAJA=1
 
 
 EXTRA_LINK_FLAGS = -L $(RAJA_LOCATION)/lib -lamdhip64
diff --git a/src/cuda_foralls.h b/src/cuda_foralls.h
@@ -202,6 +202,12 @@ void forall3(T1 &irange, T2 &jrange, T3 &krange, LoopBody &&body) {
   cudaStreamSynchronize(0);
 }
 
+template <typename T1, typename T2, typename T3, typename LoopBody>
+void forall3asyncnotimer(T1 &irange, T2 &jrange, T3 &krange, LoopBody &&body) {
+  forall3async(irange, jrange, krange, body);
+}
+  
+
 template <typename T1, typename T2, typename T3, typename LoopBody>
 void forall3GSasync(T1 &irange, T2 &jrange, T3 &krange, LoopBody &&body) {
   dim3 tpb(irange.tpb, jrange.tpb, krange.tpb);
diff --git a/src/curvitest.C b/src/curvitest.C
@@ -18,6 +18,7 @@ class Sarray {
   Sarray(int nc, int ibeg, int iend, int jbeg, int jend, int kbeg, int kend);
   std::string fill(std::istringstream& iss);
   void init();
+  void init2();
   double norm();
   int m_nc, m_ni, m_nj, m_nk;
   int m_ib, m_ie, m_jb, m_je, m_kb, m_ke;
@@ -26,6 +27,7 @@ class Sarray {
   double* m_data;
   size_t size;
   int g;
+  
 };
 
 std::string Sarray::fill(std::istringstream& iss) {
@@ -66,6 +68,37 @@ void Sarray::init() {
   forallasync(0, size / 8,
               [=] __device__(int i) { lm_data[i] = sin(double(i)); });
 }
+
+void Sarray::init2() {
+
+  Range<64> I(0,m_ni);
+  Range<2> J(0,m_nj);
+  Range<2> K(0,m_nk); 
+
+  double dx = 0.001;
+    int nc = m_nc;
+  int offi = nc;
+  int offj = nc*m_ni;
+  int offk = nc*m_ni*m_nj;
+
+  double *data = m_data;
+  
+ 
+  forall3asyncnotimer(
+	       I, J, K, [=] __device__(int i, int j, int k) {
+		 for (int c=0;c<nc;c++){
+		   
+		   int indx = c + i * offi + j * offj +
+		     k * offk;
+		   double x = i*dx;
+		   double y = j*dx;
+		   double z = k*dx;
+		   double f = sin(x)*sin(y)*sin(z);
+		   data[indx]=f;
+	       }
+	       });
+  
+}
 double Sarray::norm() {
   double ret = 0.0;
   for (size_t i = 0; i < size / 8; i++) ret += m_data[i] * m_data[i];
@@ -125,19 +158,21 @@ int main(int argc, char* argv[]) {
       std::cout << x.first << " " << x.second->g << " " << x.second->m_npts
                 << "\n";
 #endif
-      x.second->init();
+      x.second->init2();
     }
-#ifdef VERBOSE
-  std::cout << "Done with map array output\n";
-#endif
+
 
 #ifdef ENABLE_CUDA
   cudaStreamSynchronize(0);
-#endif
+ #endif
 #ifdef ENABLE_HIP
   hipStreamSynchronize(0);
 #endif
 
+#ifdef VERBOSE
+  std::cout << "Done with map array output\n";
+#endif
+
   void* ptr;
 #ifdef ENABLE_CUDA
   if (cudaMallocManaged(&ptr, (6 + 384 + 24 + 48 + 6 + 384 + 6 + 6) *
@@ -222,7 +257,8 @@ int main(int argc, char* argv[]) {
               << arrays[i]["a_Uacc"]->norm() << "\n";
     std::cout << "Norm of output " << std::defaultfloat << std::setprecision(20)
               << arrays[i]["a_Uacc"]->norm() << "\n";
-    const double exact_norm = 9.86238393426104e+17;
+    //const double exact_norm = 9.86238393426104e+17;
+    const double exact_norm = 202.0512747393526638; // for init2
     double err = (arrays[i]["a_Uacc"]->norm() - exact_norm) / exact_norm * 100;
     std::cout << "Error = " << std::setprecision(2) << err << " %\n";
   }
diff --git a/src/hip_foralls.h b/src/hip_foralls.h
@@ -164,11 +164,26 @@ void forall3(int start0, int end0, int start1, int end1, int start2, int end2,
   printf("Launching the kernel 3d \n");
   // forall3kernel<<<blocks, tpb>>>(start0, end0, start1, end1, start2, end2,
   // body);
-  hipLaunchKernelGGL(forall3kernel, blocks, tpb, 0, 0, start0, end0, start1,
+  hipLaunchKernelGGL(forall3kernel, blocks, tpb, 0, 0, 0,start0, end0, start1,
                      end1, start2, end2, body);
   hipDeviceSynchronize();
 }
 
+template <typename T1, typename T2, typename T3, typename LoopBody>
+void forall3asyncnotimer(T1 &irange, T2 &jrange, T3 &krange, LoopBody &&body) {
+  dim3 tpb(irange.tpb, jrange.tpb, krange.tpb);
+  dim3 blocks(irange.blocks, jrange.blocks, krange.blocks);
+
+  // forall3kernel<<<blocks, tpb>>>(irange.start, irange.end, jrange.start,
+  //                              jrange.end, krange.start, krange.end, body);
+  //hipEvent_t start, stop;
+  //hipEventCreate(&start);
+  //hipEventCreate(&stop);
+  hipLaunchKernelGGL(forall3kernel, blocks, tpb, 0, 0,
+                        irange.start, irange.end, jrange.start, jrange.end,
+                        krange.start, krange.end, body);
+}
+
 template <typename T1, typename T2, typename T3, typename LoopBody>
 void forall3async(T1 &irange, T2 &jrange, T3 &krange, LoopBody &&body) {
   dim3 tpb(irange.tpb, jrange.tpb, krange.tpb);
