Added Jacobi example using OpenACC with Unified Memory

Author: nsakharnykh

posts/openacc-unified-memory-jacobi/Makefile

@@ -0,0 +1,33 @@
+#
+#  Copyright 2015 NVIDIA Corporation
+#
+#  Licensed under the Apache License, Version 2.0 (the "License");
+#  you may not use this file except in compliance with the License.
+#  You may obtain a copy of the License at
+#
+#      http://www.apache.org/licenses/LICENSE-2.0
+#
+#  Unless required by applicable law or agreed to in writing, software
+#  distributed under the License is distributed on an "AS IS" BASIS,
+#  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+#  See the License for the specific language governing permissions and
+#  limitations under the License.
+#
+
+CC       = pgcc
+CCFLAGS  = -I./common
+ACCFLAGS = -acc -ta=tesla:managed -Minfo=accel
+OMPFLAGS = -fast -mp -Minfo
+
+BIN =  laplace2d_omp laplace2d_acc
+
+all: $(BIN)
+
+laplace2d_acc: laplace2d.c
+	$(CC) $(CCFLAGS) $(ACCFLAGS) -o $@ $<
+
+laplace2d_omp: laplace2d.c
+	$(CC) $(CCFLAGS) $(OMPFLAGS) -o $@ $<
+
+clean:
+	$(RM) $(BIN)

posts/openacc-unified-memory-jacobi/common/timer.h

@@ -0,0 +1,78 @@
+/* Copyright (c) 2012, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+ 
+#ifndef TIMER_H
+#define TIMER_H
+
+#include <stdlib.h>
+
+#ifdef WIN32
+#define WIN32_LEAN_AND_MEAN
+#include <windows.h>
+#else
+#include <sys/time.h>
+#endif
+
+#ifdef WIN32
+double PCFreq = 0.0;
+__int64 timerStart = 0;
+#else
+struct timeval timerStart;
+#endif
+
+void StartTimer()
+{
+#ifdef WIN32
+    LARGE_INTEGER li;
+    if(!QueryPerformanceFrequency(&li))
+        printf("QueryPerformanceFrequency failed!\n");
+
+    PCFreq = (double)li.QuadPart/1000.0;
+
+    QueryPerformanceCounter(&li);
+    timerStart = li.QuadPart;
+#else
+    gettimeofday(&timerStart, NULL);
+#endif
+}
+
+// time elapsed in ms
+double GetTimer()
+{
+#ifdef WIN32
+    LARGE_INTEGER li;
+    QueryPerformanceCounter(&li);
+    return (double)(li.QuadPart-timerStart)/PCFreq;
+#else
+    struct timeval timerStop, timerElapsed;
+    gettimeofday(&timerStop, NULL);
+    timersub(&timerStop, &timerStart, &timerElapsed);
+    return timerElapsed.tv_sec*1000.0+timerElapsed.tv_usec/1000.0;
+#endif
+}
+
+#endif // TIMER_H

posts/openacc-unified-memory-jacobi/laplace2d.c

@@ -0,0 +1,101 @@
+/*
+ *  Copyright 2015 NVIDIA Corporation
+ *
+ *  Licensed under the Apache License, Version 2.0 (the "License");
+ *  you may not use this file except in compliance with the License.
+ *  You may obtain a copy of the License at
+ *
+ *      http://www.apache.org/licenses/LICENSE-2.0
+ *
+ *  Unless required by applicable law or agreed to in writing, software
+ *  distributed under the License is distributed on an "AS IS" BASIS,
+ *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ *  See the License for the specific language governing permissions and
+ *  limitations under the License.
+ */
+
+#include <math.h>
+#include <string.h>
+#include "timer.h"
+
+int main(int argc, char** argv)
+{
+    int n = 4096;
+    int m = 4096;
+    int iter_max = 1000;
+    
+    const float pi  = 2.0f * asinf(1.0f);
+    const float tol = 1.0e-5f;
+    float error     = 1.0f;
+    
+    float *restrict A = (float*)malloc(sizeof(float)*n*m);
+    float *restrict Anew = (float*)malloc(sizeof(float)*n*m);
+    float *restrict y0 = (float*)malloc(sizeof(float)*n);
+
+    memset(A, 0, n * m * sizeof(float));
+    
+    // set boundary conditions
+    for (int i = 0; i < m; i++)
+    {
+        A[0*m+i]   = 0.f;
+        A[(n-1)*m+i] = 0.f;
+    }
+    
+    for (int j = 0; j < n; j++)
+    {
+        y0[j] = sinf(pi * j / (n-1));
+        A[j*m+0] = y0[j];
+        A[j*m+m-1] = y0[j]*expf(-pi);
+    }
+    
+    printf("Jacobi relaxation Calculation: %d x %d mesh\n", n, m);
+    
+    StartTimer();
+    int iter = 0;
+    
+#pragma omp parallel for shared(Anew)
+    for (int i = 1; i < m; i++)
+    {
+       Anew[0*m+i]   = 0.f;
+       Anew[(n-1)*m+i] = 0.f;
+    }
+#pragma omp parallel for shared(Anew)    
+    for (int j = 1; j < n; j++)
+    {
+        Anew[j*m+0]   = y0[j];
+        Anew[j*m+m-1] = y0[j]*expf(-pi);
+    }
+    
+    while ( error > tol && iter < iter_max ) {
+        error = 0.f;
+
+    #pragma acc kernels 
+    {
+        #pragma acc loop independent
+        for( int j = 1; j < n-1; j++) {
+            for( int i = 1; i < m-1; i++ ) {
+
+                Anew[j*m+i] = 0.25f * ( A[j*m+i+1] + A[j*m+i-1]
+                                     + A[(j-1)*m+i] + A[(j+1)*m+i]);
+
+                error = fmaxf( error, fabsf(Anew[j*m+i]-A[j*m+i]));
+            }
+        }
+
+        #pragma acc loop independent
+        for( int j = 1; j < n-1; j++) {
+            for( int i = 1; i < m-1; i++ ) {
+                A[j*m+i] = Anew[j*m+i];    
+            }
+        }
+    }
+
+        if(iter % 100 == 0) printf("%5d, %0.6f\n", iter, error);
+        
+        iter++;
+    }
+
+    double runtime = GetTimer();
+ 
+    printf(" total: %f s\n", runtime / 1000.f);
+}