Markdown fixes & Whitespace cleanup for samples (#1096)
* Fix multiline code blocks in README's
* Whitespace cleanup
[ROCm/clr commit: fb92feae0e]
This commit is contained in:
Nick Curtis
committad av
Maneesh Gupta
Maneesh Gupta
förälder
89da742110
incheckning
3b6b356d23
@@ -7,30 +7,32 @@ This tutorial shows how to get write simple HIP application. We will write the s
|
||||
HIP is a C++ runtime API and kernel language that allows developers to create portable applications that can run on AMD and other GPU’s. Our goal was to rise above the lowest-common-denominator paths and deliver a solution that allows you, the developer, to use essential hardware features and maximize your application’s performance on GPU hardware.
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
Here is simple example showing how to write your first program in HIP.
|
||||
In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu).
|
||||
In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu).
|
||||
|
||||
## Device-side code
|
||||
We will work on device side code first, Here is simple example showing a snippet of HIP device side code:
|
||||
|
||||
`__global__ void matrixTranspose(float *out, `
|
||||
` float *in, `
|
||||
` const int width, `
|
||||
` const int height) `
|
||||
`{ `
|
||||
` int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; `
|
||||
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; `
|
||||
` `
|
||||
` out[y * width + x] = in[x * height + y]; `
|
||||
`} `
|
||||
```
|
||||
__global__ void matrixTranspose(float *out,
|
||||
float *in,
|
||||
const int width,
|
||||
const int height)
|
||||
{
|
||||
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
out[y * width + x] = in[x * height + y];
|
||||
}
|
||||
```
|
||||
|
||||
`__global__` keyword is the Function-Type Qualifiers, it is used with functions that are executed on device and are called/launched from the hosts.
|
||||
other function-type qualifiers are:
|
||||
@@ -40,11 +42,11 @@ other function-type qualifiers are:
|
||||
`__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "hipThreadIdx_x", will talk about it latter). A possible workaround is to pass the necessary coordinate info as an argument to the function.
|
||||
`__host__` cannot combine with `__global__`.
|
||||
|
||||
`__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*.
|
||||
`__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*.
|
||||
|
||||
Next keyword is `void`. HIP `__global__` functions must have a `void` return type. Global functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute.
|
||||
|
||||
The kernel function begins with
|
||||
The kernel function begins with
|
||||
` int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;`
|
||||
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;`
|
||||
here the keyword hipBlockIdx_x, hipBlockIdx_y and hipBlockIdx_z(not used here) are the built-in functions to identify the threads in a block. The keyword hipBlockDim_x, hipBlockDim_y and hipBlockDim_z(not used here) are to identify the dimensions of the block.
|
||||
@@ -60,18 +62,20 @@ We allocated memory to the Matrix on host side by using malloc and initiallized
|
||||
here the first parameter is the destination pointer, second is the source pointer, third is the size of memory copy and the last specify the direction on memory copy(which is in this case froom host to device). While in order to transfer memory from device to host, use `hipMemcpyDeviceToHost` and for device to device memory copy use `hipMemcpyDeviceToDevice`.
|
||||
|
||||
Now, we'll see how to launch the kernel.
|
||||
` hipLaunchKernelGGL(matrixTranspose, `
|
||||
` dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), `
|
||||
` dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), `
|
||||
` 0, 0, `
|
||||
` gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT); `
|
||||
```
|
||||
hipLaunchKernelGGL(matrixTranspose,
|
||||
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
|
||||
0, 0,
|
||||
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
|
||||
```
|
||||
|
||||
HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the `Cuda <<< >>>` syntax). In HIP,
|
||||
- Kernels launch with the `"hipLaunchKernelGGL"` function
|
||||
- The first five parameters to hipLaunchKernelGGL are the following:
|
||||
- **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. In current application it's "matrixTranspose".
|
||||
- **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)".
|
||||
- **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)".
|
||||
- **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)".
|
||||
- **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)".
|
||||
- **size_t dynamicShared**: amount of additional shared memory to allocate when launching the kernel. In MatrixTranspose sample, it's '0'.
|
||||
- **hipStream_t**: stream where the kernel should execute. A value of 0 corresponds to the NULL stream.In MatrixTranspose sample, it's '0'.
|
||||
- Kernel arguments follow these first five parameters. Here, these are "gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT".
|
||||
|
||||
@@ -6,7 +6,7 @@ This tutorial is about how to use inline GCN asm in kernel. In this tutorial, we
|
||||
|
||||
If you want to take advantage of the extra performance benefits of writing in assembly as well as take advantage of special GPU hardware features that were only available through assemby, then this tutorial is for you. In this tutorial we'll be explaining how to start writing inline asm in kernel.
|
||||
|
||||
For more insight Please read the following blogs by Ben Sander
|
||||
For more insight Please read the following blogs by Ben Sander
|
||||
[The Art of AMDGCN Assembly: How to Bend the Machine to Your Will](gpuopen.com/amdgcn-assembly)
|
||||
[AMD GCN Assembly: Cross-Lane Operations](http://gpuopen.com/amd-gcn-assembly-cross-lane-operations/)
|
||||
|
||||
@@ -15,13 +15,13 @@ For more information:
|
||||
[User Guide for AMDGPU Back-end](llvm.org/docs/AMDGPUUsage.html)
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the our very first tutorial.
|
||||
|
||||
|
||||
@@ -7,13 +7,13 @@ This tutorial is follow-up of the previous one where we learn how to write our f
|
||||
Memory transfer and kernel execution are the most important parameter in parallel computing (specially HPC and machine learning). Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore obtaining the memory transfer timing and kernel execution timing plays key role in application optimization.
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to get the performance score for memory transfer and kernel execution time.
|
||||
|
||||
@@ -21,12 +21,16 @@ We will be using the Simple Matrix Transpose application from the previous tutor
|
||||
|
||||
We'll learn how to use the event management functionality of HIP runtime api. In the same sourcecode, we used for MatrixTranspose we will declare the following events as follows:
|
||||
|
||||
` hipEvent_t start, stop;`
|
||||
```
|
||||
hipEvent_t start, stop;
|
||||
```
|
||||
|
||||
We'll create the event with the help of following code:
|
||||
|
||||
` hipEventCreate(&start);`
|
||||
` hipEventCreate(&stop);`
|
||||
```
|
||||
hipEventCreate(&start);
|
||||
hipEventCreate(&stop);
|
||||
```
|
||||
|
||||
We'll use the "eventMs" variable to store the time taken value:
|
||||
` float eventMs = 1.0f;`
|
||||
@@ -41,11 +45,13 @@ Now, we'll have the operation for which we need to compute the time taken. For t
|
||||
` hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);`
|
||||
|
||||
and for kernel execution time we'll use `hipKernelLaunch`:
|
||||
` hipLaunchKernelGGL(matrixTranspose, `
|
||||
` dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), `
|
||||
` dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), `
|
||||
` 0, 0, `
|
||||
` gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT); `
|
||||
```
|
||||
hipLaunchKernelGGL(matrixTranspose,
|
||||
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
|
||||
0, 0,
|
||||
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
|
||||
```
|
||||
|
||||
Now to mark the end of the eventRecord, we will again use the hipEventRecord by passing the stop event:
|
||||
` hipEventRecord(stop, NULL);`
|
||||
|
||||
@@ -7,13 +7,13 @@ Earlier we learned how to write our first hip program, in which we compute Matri
|
||||
As we mentioned earlier that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use static shared memory and will explain the dynamic one latter.
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
|
||||
|
||||
|
||||
@@ -5,23 +5,25 @@ In this tutorial, we'll explain how to use the warp shfl operations to improve t
|
||||
## Introduction:
|
||||
|
||||
Let's talk about Warp first. The kernel code is executed in groups of fixed number of threads known as Warp. For nvidia WarpSize is 32 while for AMD, 32 for Polaris architecture and 64 for rest. Threads in a warp are referred to as lanes and are numbered from 0 to warpSize -1. With the help of shfl ops, we can directly exchange values of variable between threads without using any memory ops within a warp. There are four types of shfl ops:
|
||||
` int __shfl (int var, int srcLane, int width=warpSize); `
|
||||
` float __shfl (float var, int srcLane, int width=warpSize); `
|
||||
` int __shfl_up (int var, unsigned int delta, int width=warpSize); `
|
||||
` float __shfl_up (float var, unsigned int delta, int width=warpSize); `
|
||||
` int __shfl_down (int var, unsigned int delta, int width=warpSize); `
|
||||
` float __shfl_down (float var, unsigned int delta, int width=warpSize); `
|
||||
` int __shfl_xor (int var, int laneMask, int width=warpSize) `
|
||||
` float __shfl_xor (float var, int laneMask, int width=warpSize); `
|
||||
```
|
||||
int __shfl (int var, int srcLane, int width=warpSize);
|
||||
float __shfl (float var, int srcLane, int width=warpSize);
|
||||
int __shfl_up (int var, unsigned int delta, int width=warpSize);
|
||||
float __shfl_up (float var, unsigned int delta, int width=warpSize);
|
||||
int __shfl_down (int var, unsigned int delta, int width=warpSize);
|
||||
float __shfl_down (float var, unsigned int delta, int width=warpSize);
|
||||
int __shfl_xor (int var, int laneMask, int width=warpSize)
|
||||
float __shfl_xor (float var, int laneMask, int width=warpSize);
|
||||
```
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
|
||||
|
||||
|
||||
@@ -5,31 +5,35 @@ This tutorial is follow-up of the previous tutorial, where we learned how to use
|
||||
## Introduction:
|
||||
|
||||
Let's talk about Warp first. The kernel code is executed in groups of fixed number of threads known as Warp. For nvidia WarpSize is 32 while for AMD, 32 for Polaris architecture and 64 for rest. Threads in a warp are referred to as lanes and are numbered from 0 to warpSize -1. With the help of shfl ops, we can directly exchange values of variable between threads without using any memory ops within a warp. There are four types of shfl ops:
|
||||
` int __shfl (int var, int srcLane, int width=warpSize); `
|
||||
` float __shfl (float var, int srcLane, int width=warpSize); `
|
||||
` int __shfl_up (int var, unsigned int delta, int width=warpSize); `
|
||||
` float __shfl_up (float var, unsigned int delta, int width=warpSize); `
|
||||
` int __shfl_down (int var, unsigned int delta, int width=warpSize); `
|
||||
` float __shfl_down (float var, unsigned int delta, int width=warpSize); `
|
||||
` int __shfl_xor (int var, int laneMask, int width=warpSize) `
|
||||
` float __shfl_xor (float var, int laneMask, int width=warpSize); `
|
||||
```
|
||||
int __shfl (int var, int srcLane, int width=warpSize);
|
||||
float __shfl (float var, int srcLane, int width=warpSize);
|
||||
int __shfl_up (int var, unsigned int delta, int width=warpSize);
|
||||
float __shfl_up (float var, unsigned int delta, int width=warpSize);
|
||||
int __shfl_down (int var, unsigned int delta, int width=warpSize);
|
||||
float __shfl_down (float var, unsigned int delta, int width=warpSize);
|
||||
int __shfl_xor (int var, int laneMask, int width=warpSize);
|
||||
float __shfl_xor (float var, int laneMask, int width=warpSize);
|
||||
```
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
|
||||
|
||||
## __shfl ops in 2D
|
||||
|
||||
In the same sourcecode, we used for MatrixTranspose. We'll add the following:
|
||||
` int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; `
|
||||
` out[x*width + y] = __shfl(val,y*width + x); `
|
||||
```
|
||||
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
out[x*width + y] = __shfl(val,y*width + x);
|
||||
```
|
||||
|
||||
With the help of this application, we can say that kernel code can be converted into multi-dimensional threads with ease.
|
||||
|
||||
|
||||
@@ -7,13 +7,13 @@ Earlier we learned how to use static shared memory. In this tutorial, we'll expl
|
||||
As we mentioned earlier that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use dynamic shared memory.
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
|
||||
|
||||
@@ -25,11 +25,13 @@ Shared memory is way more faster than that of global and constant memory and acc
|
||||
here the first parameter is the data type while the second one is the variable name.
|
||||
|
||||
The other important change is:
|
||||
` hipLaunchKernelGGL(matrixTranspose, `
|
||||
```
|
||||
hipLaunchKernelGGL(matrixTranspose,
|
||||
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
|
||||
sizeof(float)*WIDTH*WIDTH, 0,
|
||||
gpuTransposeMatrix , gpuMatrix, WIDTH);
|
||||
```
|
||||
here we replaced 4th parameter with amount of additional shared memory to allocate when launching the kernel.
|
||||
|
||||
## How to build and run:
|
||||
|
||||
@@ -7,13 +7,13 @@ In all Earlier tutorial we used single stream, In this tutorial, we'll explain h
|
||||
The various instances of kernel to be executed on device in exact launch order defined by Host are called streams. We can launch multiple streams on a single device. We will learn how to learn two streams which can we scaled with ease.
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to launch multiple streams.
|
||||
|
||||
@@ -23,22 +23,28 @@ In this tutorial, we'll use both instances of shared memory (i.e., static and dy
|
||||
` hipStream_t streams[num_streams]; `
|
||||
|
||||
and create stream using `hipStreamCreate` as follows:
|
||||
` for(int i=0;i<num_streams;i++) `
|
||||
` hipStreamCreate(&streams[i]); `
|
||||
```
|
||||
for(int i=0;i<num_streams;i++)
|
||||
hipStreamCreate(&streams[i]);
|
||||
```
|
||||
|
||||
and while kernel launch, we make the following changes in 5th parameter to hipLaunchKernelGGL(having 0 as the default stream value):
|
||||
|
||||
` hipLaunchKernelGGL(matrixTranspose_static_shared, `
|
||||
```
|
||||
hipLaunchKernelGGL(matrixTranspose_static_shared,
|
||||
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
|
||||
0, streams[0],
|
||||
gpuTransposeMatrix[0], data[0], width);
|
||||
```
|
||||
|
||||
` hipLaunchKernelGGL(matrixTranspose_dynamic_shared, `
|
||||
```
|
||||
hipLaunchKernelGGL(matrixTranspose_dynamic_shared,
|
||||
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
|
||||
sizeof(float)*WIDTH*WIDTH, streams[1],
|
||||
gpuTransposeMatrix[1], data[1], width);
|
||||
```
|
||||
|
||||
here we replaced 4th parameter with amount of additional shared memory to allocate when launching the kernel.
|
||||
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
## Using Pragma unroll ###
|
||||
|
||||
In this tutorial, we'll explain how to use #pragma unroll to improve the performance.
|
||||
In this tutorial, we'll explain how to use #pragma unroll to improve the performance.
|
||||
|
||||
## Introduction:
|
||||
|
||||
@@ -8,24 +8,26 @@ Loop unrolling optimization hints can be specified with #pragma unroll and #prag
|
||||
Specifying #pragma unroll without a parameter directs the loop unroller to attempt to fully unroll the loop if the trip count is known at compile time and attempt to partially unroll the loop if the trip count is not known at compile time.
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
|
||||
|
||||
## prerequiste knowledge:
|
||||
|
||||
Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
|
||||
|
||||
## Simple Matrix Transpose
|
||||
## Simple Matrix Transpose
|
||||
|
||||
For this tutorial we will be using MatrixTranspose with shfl operation i.e., our 4_shfl tutorial since it is the only examples where we used loops inside the kernel.
|
||||
|
||||
In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for MatrixTranspose. We'll add it just before the for loop as following:
|
||||
|
||||
`#pragma unroll `
|
||||
` for(int i=0;i<width;i++) `
|
||||
` { `
|
||||
` for(int j=0;j<width;j++) `
|
||||
` out[i*width + j] = __shfl(val,j*width + i); `
|
||||
` } `
|
||||
```
|
||||
#pragma unroll
|
||||
for(int i=0;i<width;i++)
|
||||
{
|
||||
for(int j=0;j<width;j++)
|
||||
out[i*width + j] = __shfl(val,j*width + i);
|
||||
}
|
||||
```
|
||||
|
||||
Specifying the optional parameter, #pragma unroll value, directs the unroller to unroll the loop value times. Be careful while using it.
|
||||
Specifying #pragma nounroll indicates that the loop should not be unroll. #pragma unroll 1 will show the same behaviour.
|
||||
|
||||
Referens i nytt ärende
Block a user