Update example source files

Author: neon60

docs/tools/example_codes/device_selection.hip

@@ -40,10 +40,11 @@
     }                                                \
 }
 
-__global__ void simpleKernel(double *data)
+__global__ void simpleKernel(double *data, std::size_t elems)
 {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    data[idx] = idx * 2.0;
+    if(idx < elems)
+        data[idx] = idx * 2.0;
 }
 
 int main()
@@ -58,30 +59,31 @@ int main()
 
     double* deviceData0;
     double* deviceData1;
-    std::size_t size = 1024 * sizeof(*deviceData0);
+    constexpr std::size_t elems = 1024;
+    constexpr std::size_t size = elems * sizeof(double);
 
     int deviceId0 = 0;
     int deviceId1 = 1;
 
     // Set device 0 and perform operations
     HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
     HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
-    simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
+    simpleKernel<<<8, 128>>>(deviceData0, elems); // Launch kernel on device 0
     HIP_CHECK(hipDeviceSynchronize());
 
     // Set device 1 and perform operations
     HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
     HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
-    simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
+    simpleKernel<<<8, 128>>>(deviceData1, elems); // Launch kernel on device 1
     HIP_CHECK(hipDeviceSynchronize());
 
     // Copy result from device 0
-    double hostData0[1024];
+    double hostData0[elems];
     HIP_CHECK(hipSetDevice(deviceId0));
     HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));
 
     // Copy result from device 1
-    double hostData1[1024];
+    double hostData1[elems];
     HIP_CHECK(hipSetDevice(deviceId1));
     HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));
 

docs/tools/example_codes/multi_device_synchronization.hip

@@ -40,10 +40,11 @@
     }                                                \
 }
 
-__global__ void simpleKernel(double *data)
+__global__ void simpleKernel(double *data, std::size_t elems)
 {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    data[idx] = idx * 2.0;
+    if(idx < elems)
+        data[idx] = idx * 2.0;
 }
 
 int main()
@@ -57,8 +58,10 @@ int main()
         return EXIT_SUCCESS;
     }
 
-    double *deviceData0, *deviceData1;
-    std::size_t size = 1024 * sizeof(*deviceData0);
+    double *deviceData0;
+    double *deviceData1;
+    constexpr std::size_t elems = 1024;
+    constexpr std::size_t size = elems * sizeof(double);
 
     // Create streams and events for each device
     hipStream_t stream0, stream1;
@@ -83,7 +86,7 @@ int main()
     HIP_CHECK(hipEventRecord(startEvent0, stream0));
 
     // Launch the kernel asynchronously on device 0
-    simpleKernel<<<1000, 128, 0, stream0>>>(deviceData0);
+    simpleKernel<<<8, 128, 0, stream0>>>(deviceData0, elems);
 
     // Record the stop event on device 0
     HIP_CHECK(hipEventRecord(stopEvent0, stream0));
@@ -96,7 +99,7 @@ int main()
     HIP_CHECK(hipEventRecord(startEvent1, stream1));
 
     // Launch the kernel asynchronously on device 1
-    simpleKernel<<<1000, 128, 0, stream1>>>(deviceData1);
+    simpleKernel<<<8, 128, 0, stream1>>>(deviceData1, elems);
 
     // Record the stop event on device 1
     HIP_CHECK(hipEventRecord(stopEvent1, stream1));

docs/tools/example_codes/p2p_memory_access.hip

@@ -40,10 +40,11 @@
     }                                                \
 }
 
-__global__ void simpleKernel(double *data)
+__global__ void simpleKernel(double *data, std::size_t elems)
 {
-    int idx   = blockIdx.x * blockDim.x + threadIdx.x;
-    data[idx] = idx * 2.0;
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if(idx < elems)
+        data[idx] = idx * 2.0;
 }
 
 int main()
@@ -58,7 +59,8 @@ int main()
 
     double* deviceData0;
     double* deviceData1;
-    std::size_t size = 1024 * sizeof(*deviceData0);
+    constexpr std::size_t elems = 1024;
+    constexpr std::size_t size = elems * sizeof(double);
 
     int deviceId0 = 0;
     int deviceId1 = 1;
@@ -74,13 +76,13 @@ int main()
     // Set device 0 and perform operations
     HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
     HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
-    simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
+    simpleKernel<<<8, 128>>>(deviceData0, elems); // Launch kernel on device 0
     HIP_CHECK(hipDeviceSynchronize());
 
     // Set device 1 and perform operations
     HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
     HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
-    simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
+    simpleKernel<<<8, 128>>>(deviceData1, elems); // Launch kernel on device 1
     HIP_CHECK(hipDeviceSynchronize());
 
     // Use peer-to-peer access
@@ -90,12 +92,12 @@ int main()
     HIP_CHECK(hipMemcpy(deviceData0, deviceData1, size, hipMemcpyDeviceToDevice));
 
     // Copy result from device 0
-    double hostData0[1024];
+    double hostData0[elems];
     HIP_CHECK(hipSetDevice(deviceId0));
     HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));
 
     // Copy result from device 1
-    double hostData1[1024];
+    double hostData1[elems];
     HIP_CHECK(hipSetDevice(deviceId1));
     HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));
 

docs/tools/example_codes/p2p_memory_access_host_staging.hip

@@ -40,10 +40,11 @@
     }                                                \
 }
 
-__global__ void simpleKernel(double *data)
+__global__ void simpleKernel(double *data, std::size_t elems)
 {
-    int idx   = blockIdx.x * blockDim.x + threadIdx.x;
-    data[idx] = idx * 2.0;
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if(idx < elems)
+        data[idx] = idx * 2.0;
 }
 
 int main()
@@ -58,34 +59,35 @@ int main()
 
     double* deviceData0;
     double* deviceData1;
-    std::size_t size = 1024 * sizeof(*deviceData0);
+    constexpr std::size_t elems = 1024;
+    constexpr std::size_t size = elems * sizeof(double);
 
     int deviceId0 = 0;
     int deviceId1 = 1;
 
     // Set device 0 and perform operations
     HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
     HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
-    simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
+    simpleKernel<<<8, 128>>>(deviceData0, elems); // Launch kernel on device 0
     HIP_CHECK(hipDeviceSynchronize());
 
     // Set device 1 and perform operations
     HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
     HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
-    simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
+    simpleKernel<<<8, 128>>>(deviceData1, elems); // Launch kernel on device 1
     HIP_CHECK(hipDeviceSynchronize());
 
     // Use deviceData0 on device 1. This works but incurs a performance penalty.
     HIP_CHECK(hipSetDevice(deviceId1));
     HIP_CHECK(hipMemcpy(deviceData1, deviceData0, size, hipMemcpyDeviceToDevice));
 
     // Copy result from device 0
-    double hostData0[1024];
+    double hostData0[elems];
     HIP_CHECK(hipSetDevice(deviceId0));
     HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));
 
     // Copy result from device 1
-    double hostData1[1024];
+    double hostData1[elems];
     HIP_CHECK(hipSetDevice(deviceId1));
     HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));