update examples

examples/barrier.cpp

@@ -1,3 +1,9 @@
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <cstdlib>
 #include <functional>
 #include <iostream>

examples/cuda/multiple_streams.cu

@@ -1,17 +1,9 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 
 #include <chrono>
 #include <cstdlib>
@@ -30,14 +22,10 @@
 #include <boost/fiber/cuda/waitfor.hpp>
 
 __global__
-void kernel( int size, int * a, int * b, int * c) {
+void vector_add( int * a, int * b, int * c, int size) {
     int idx = threadIdx.x + blockIdx.x * blockDim.x;
     if ( idx < size) {
-        int idx1 = (idx + 1) % 256;
-        int idx2 = (idx + 2) % 256;
-        float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
-        float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
-        c[idx] = (as + bs) / 2;
+        c[idx] = a[idx] + b[idx];
     }
 }
 
@@ -75,8 +63,8 @@ int main() {
                         cudaMemcpyAsync( dev_a1, host_a + i + size, size * sizeof( int), cudaMemcpyHostToDevice, stream1);
                         cudaMemcpyAsync( dev_b0, host_b + i, size * sizeof( int), cudaMemcpyHostToDevice, stream0);
                         cudaMemcpyAsync( dev_b1, host_b + i + size, size * sizeof( int), cudaMemcpyHostToDevice, stream1);
-                        kernel<<< size / 256, 256, 0, stream0 >>>( size, dev_a0, dev_b0, dev_c0);
-                        kernel<<< size / 256, 256, 0, stream1 >>>( size, dev_a1, dev_b1, dev_c1);
+                        vector_add<<< size / 256, 256, 0, stream0 >>>( dev_a0, dev_b0, dev_c0, size);
+                        vector_add<<< size / 256, 256, 0, stream1 >>>( dev_a1, dev_b1, dev_c1, size);
                         cudaMemcpyAsync( host_c + i, dev_c0, size * sizeof( int), cudaMemcpyDeviceToHost, stream0);
                         cudaMemcpyAsync( host_c + i + size, dev_c1, size * sizeof( int), cudaMemcpyDeviceToHost, stream1);
                     }

examples/cuda/single_stream.cu

@@ -1,17 +1,8 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
+
+//          Copyright Oliver Kowalke 2017.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
 
 #include <chrono>
 #include <cstdlib>
@@ -30,14 +21,10 @@
 #include <boost/fiber/cuda/waitfor.hpp>
 
 __global__
-void kernel( int size, int * a, int * b, int * c) {
+void vector_add( int * a, int * b, int * c, int size) {
     int idx = threadIdx.x + blockIdx.x * blockDim.x;
     if ( idx < size) {
-        int idx1 = (idx + 1) % 256;
-        int idx2 = (idx + 2) % 256;
-        float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
-        float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
-        c[idx] = (as + bs) / 2;
+        c[idx] = a[idx] + b[idx];
     }
 }
 
@@ -68,7 +55,7 @@ int main() {
                 for ( int i = 0; i < full_size; i += size) {
                     cudaMemcpyAsync( dev_a, host_a + i, size * sizeof( int), cudaMemcpyHostToDevice, stream);
                     cudaMemcpyAsync( dev_b, host_b + i, size * sizeof( int), cudaMemcpyHostToDevice, stream);
-                    kernel<<< size / 256, 256, 0, stream >>>( size, dev_a, dev_b, dev_c);
+                    vector_add<<< size / 256, 256, 0, stream >>>( dev_a, dev_b, dev_c, size);
                     cudaMemcpyAsync( host_c + i, dev_c, size * sizeof( int), cudaMemcpyDeviceToHost, stream);
                 }
                 auto result = boost::fibers::cuda::waitfor_all( stream);

examples/future.cpp

@@ -1,3 +1,9 @@
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <cstdlib>
 #include <exception>
 #include <functional>

examples/hip/multiple_streams.cpp

@@ -1,19 +1,9 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
 
-#include "hip/hip_runtime.h"
+//          Copyright Oliver Kowalke 2017.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <chrono>
 #include <cstdlib>
 #include <iostream>
@@ -31,14 +21,10 @@
 #include <boost/fiber/hip/waitfor.hpp>
 
 __global__
-void kernel( hipLaunchParm lp, int size, int * a, int * b, int * c) {
-    int idx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+void vector_add( int * a, int * b, int * c, int size) {
+    int idx = threadIdx.x + blockIdx.x * blockDim.x;
     if ( idx < size) {
-        int idx1 = (idx + 1) % 256;
-        int idx2 = (idx + 2) % 256;
-        float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
-        float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
-        c[idx] = (as + bs) / 2;
+        c[idx] = a[idx] + b[idx];
     }
 }
 
@@ -76,8 +62,8 @@ int main() {
                         hipMemcpyAsync( dev_a1, host_a + i + size, size * sizeof( int), hipMemcpyHostToDevice, stream1);
                         hipMemcpyAsync( dev_b0, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream0);
                         hipMemcpyAsync( dev_b1, host_b + i + size, size * sizeof( int), hipMemcpyHostToDevice, stream1);
-                        hipLaunchKernel(kernel, dim3(size / 256), dim3(256), 0, stream0 ,  size, dev_a0, dev_b0, dev_c0);
-                        hipLaunchKernel(kernel, dim3(size / 256), dim3(256), 0, stream1 ,  size, dev_a1, dev_b1, dev_c1);
+                        hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream0, dev_a0, dev_b0, dev_c0, size);
+                        hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream1, dev_a1, dev_b1, dev_c1, size);
                         hipMemcpyAsync( host_c + i, dev_c0, size * sizeof( int), hipMemcpyDeviceToHost, stream0);
                         hipMemcpyAsync( host_c + i + size, dev_c1, size * sizeof( int), hipMemcpyDeviceToHost, stream1);
                     }

examples/hip/single_stream.cpp

@@ -1,19 +1,9 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
 
-#include "hip/hip_runtime.h"
+//          Copyright Oliver Kowalke 2017.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <chrono>
 #include <cstdlib>
 #include <iostream>
@@ -31,14 +21,10 @@
 #include <boost/fiber/hip/waitfor.hpp>
 
 __global__
-void kernel( hipLaunchParm lp, int size, int * a, int * b, int * c) {
-    int idx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+void vector_add( int * a, int * b, int * c, int size) {
+    int idx = threadIdx.x + blockIdx.x * blockDim.x;
     if ( idx < size) {
-        int idx1 = (idx + 1) % 256;
-        int idx2 = (idx + 2) % 256;
-        float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
-        float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
-        c[idx] = (as + bs) / 2;
+        c[idx] = a[idx] + b[idx];
     }
 }
 
@@ -69,7 +55,7 @@ int main() {
                 for ( int i = 0; i < full_size; i += size) {
                     hipMemcpyAsync( dev_a, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream);
                     hipMemcpyAsync( dev_b, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream);
-                    hipLaunchKernel(kernel, dim3(size / 256), dim3(256), 0, stream, size, dev_a, dev_b, dev_c);
+                    hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream, dev_a, dev_b, dev_c, size);
                     hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream);
                 }
                 auto result = boost::fibers::hip::waitfor_all( stream);

examples/join.cpp

@@ -1,3 +1,9 @@
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <cstdlib>
 #include <functional>
 #include <stdexcept>

examples/ping_pong.cpp

@@ -1,3 +1,9 @@
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <cstdlib>
 #include <iostream>
 #include <stdexcept>

examples/range_for.cpp

@@ -1,3 +1,9 @@
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <cstdlib>
 #include <iostream>
 #include <stdexcept>

examples/simple.cpp

@@ -1,3 +1,9 @@
+
+//          Copyright Oliver Kowalke 2013.
+// Distributed under the Boost Software License, Version 1.0.
+//    (See accompanying file LICENSE_1_0.txt or copy at
+//          http://www.boost.org/LICENSE_1_0.txt)
+
 #include <cstdlib>
 #include <iostream>
 #include <memory>