Include examples from rocm examples source

Author: neon60

docs/how-to/hip_cpp_language_extensions.rst

@@ -103,66 +103,10 @@ The kernel arguments are listed after the configuration parameters.
 
 .. code-block:: cpp
 
-  #include <hip/hip_runtime.h>
-  #include <iostream>
-
-  #define HIP_CHECK(expression)                                \
-  {                                                            \
-      const hipError_t err = expression;                       \
-      if(err != hipSuccess){                                   \
-          std::cerr << "HIP error: " << hipGetErrorString(err) \
-              << " at " << __LINE__ << "\n";                   \
-      }                                                        \
-  }
-
-  // Performs a simple initialization of an array with the thread's index variables.
-  // This function is only available in device code.
-  __device__ void init_array(float * const a, const unsigned int arraySize){
-    // globalIdx uniquely identifies a thread in a 1D launch configuration.
-    const int globalIdx = threadIdx.x + blockIdx.x * blockDim.x;
-    // Each thread initializes a single element of the array.
-    if(globalIdx < arraySize){
-      a[globalIdx] = globalIdx;
-    }
-  }
-
-  // Rounds a value up to the next multiple.
-  // This function is available in host and device code.
-  __host__ __device__ constexpr int round_up_to_nearest_multiple(int number, int multiple){
-    return (number + multiple - 1)/multiple;
-  }
-
-  __global__ void example_kernel(float * const a, const unsigned int N)
-  {
-    // Initialize array.
-    init_array(a, N);
-    // Perform additional work:
-    // - work with the array
-    // - use the array in a different kernel
-    // - ...
-  }
-
-  int main()
-  {
-    constexpr int N = 100000000; // problem size
-    constexpr int blockSize = 256; //configurable block size
-
-    //needed number of blocks for the given problem size
-    constexpr int gridSize = round_up_to_nearest_multiple(N, blockSize);
-
-    float *a;
-    // allocate memory on the GPU
-    HIP_CHECK(hipMalloc(&a, sizeof(*a) * N));
-
-    std::cout << "Launching kernel." << std::endl;
-    example_kernel<<<dim3(gridSize), dim3(blockSize), 0/*example doesn't use shared memory*/, 0/*default stream*/>>>(a, N);
-    // make sure kernel execution is finished by synchronizing. The CPU can also
-    // execute other instructions during that time
-    HIP_CHECK(hipDeviceSynchronize());
-    std::cout << "Kernel execution finished." << std::endl;
-
-    HIP_CHECK(hipFree(a));
-  }
+  .. literalinclude:: ../tools/example_codes/calling_global_functions.hip
+      :start-after: // [sphinx-start]
+      :end-before: // [sphinx-end]
+      :language: cpp
 
 Inline qualifiers
 --------------------------------------------------------------------------------
@@ -321,28 +265,10 @@ launch has to specify the needed amount of ``extern`` shared memory in the launc
 configuration. The statically allocated shared memory is allocated without this
 parameter.
 
-.. code-block:: cpp
-
-  #include <hip/hip_runtime.h>
-
-  extern __shared__ int shared_array[];
-
-  __global__ void kernel(){
-    // initialize shared memory
-    shared_array[threadIdx.x] = threadIdx.x;
-    // use shared memory - synchronize to make sure, that all threads of the
-    // block see all changes to shared memory
-    __syncthreads();
-  }
-
-  int main(){
-    //shared memory in this case depends on the configurable block size
-    constexpr int blockSize = 256;
-    constexpr int sharedMemSize = blockSize * sizeof(int);
-    constexpr int gridSize = 2;
-
-    kernel<<<dim3(gridSize), dim3(blockSize), sharedMemSize, 0>>>();
-  }
+.. literalinclude:: ../tools/example_codes/extern_shared_memory.hip
+    :start-after: // [sphinx-start]
+    :end-before: // [sphinx-end]
+    :language: cpp
 
 __managed__
 --------------------------------------------------------------------------------
@@ -735,22 +661,18 @@ with the actual frequency.
 
 The difference between the returned values represents the cycles used.
 
-.. code-block:: cpp
-
-  __global void kernel(){
-    long long int start = clock64();
-    // kernel code
-    long long int stop = clock64();
-    long long int cycles = stop - start;
-  }
+.. literalinclude:: ../tools/example_codes/timer.hip
+    :start-after: // [sphinx-kernel-start]
+    :end-before: // [sphinx-kernel-end]
+    :language: cpp
 
 ``long long int wall_clock64()`` returns the wall clock time on the device, with a constant, fixed frequency.
 The frequency is device dependent and can be queried using:
 
-.. code-block:: cpp
-
-  int wallClkRate = 0; //in kilohertz
-  hipDeviceGetAttribute(&wallClkRate, hipDeviceAttributeWallClockRate, deviceId);
+.. literalinclude:: ../tools/example_codes/timer.hip
+    :start-after: // [sphinx-query-start]
+    :end-before: // [sphinx-query-end]
+    :language: cpp
 
 .. _atomic functions: