Merge branch 'roc-1.6.x' into master
Change-Id: I367a3940a0a9e5658abc28a7dc2bfb9cf4167dc8
[ROCm/hip-tests commit: 465bc42a12]
This commit is contained in:
Maneesh Gupta
@@ -16,13 +16,15 @@ int p_iterations = 10;
|
||||
int p_beatsperiteration=1;
|
||||
int p_device = 0;
|
||||
int p_detailed = 0;
|
||||
bool p_async = 0;
|
||||
bool p_async = 0;
|
||||
int p_alignedhost = 0; // align host allocs to this granularity, in bytes. 64 or 4096 are good values to try.
|
||||
int p_onesize = 0;
|
||||
int p_onesize = 0;
|
||||
|
||||
bool p_h2d = true;
|
||||
bool p_d2h = true;
|
||||
bool p_bidir = true;
|
||||
bool p_p2p = false;
|
||||
|
||||
|
||||
//#define NO_CHECK
|
||||
|
||||
@@ -70,7 +72,7 @@ std::string sizeToString(int size)
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind)
|
||||
hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind )
|
||||
{
|
||||
if (p_async) {
|
||||
return hipMemcpyAsync(dst, src, sizeBytes, kind, NULL);
|
||||
@@ -632,6 +634,9 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB)
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
#define failed(...) \
|
||||
printf ("error: ");\
|
||||
printf (__VA_ARGS__);\
|
||||
@@ -646,6 +651,326 @@ int parseInt(const char *str, int *output)
|
||||
}
|
||||
|
||||
|
||||
void checkPeer2PeerSupport()
|
||||
{
|
||||
int deviceCnt;
|
||||
hipGetDeviceCount(&deviceCnt);
|
||||
std::cout << "Total no. of available gpu #" << deviceCnt << "\n" << std::endl;
|
||||
|
||||
for(int deviceId=0; deviceId<deviceCnt; deviceId++)
|
||||
{
|
||||
hipDeviceProp_t props;
|
||||
hipGetDeviceProperties(&props, deviceId);
|
||||
std::cout << "for gpu#" << deviceId << " " << props.name << std::endl;
|
||||
std::cout << " peer2peer supported : ";
|
||||
int PeerCnt=0;
|
||||
for (int i=0; i<deviceCnt; i++) {
|
||||
int isPeer;
|
||||
hipDeviceCanAccessPeer(&isPeer, i, deviceId);
|
||||
if (isPeer) {
|
||||
std::cout << "gpu#" << i << " ";
|
||||
++PeerCnt;
|
||||
}
|
||||
}
|
||||
if (PeerCnt==0)
|
||||
std::cout << "NONE" << " ";
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << " peer2peer not supported : ";
|
||||
int nonPeerCnt=0;
|
||||
for (int i=0; i<deviceCnt; i++) {
|
||||
int isPeer;
|
||||
hipDeviceCanAccessPeer(&isPeer, i, deviceId);
|
||||
if (!isPeer && (i!=deviceId)) {
|
||||
std::cout << "gpu#" << i << " ";
|
||||
++nonPeerCnt;
|
||||
}
|
||||
}
|
||||
if (nonPeerCnt==0)
|
||||
std::cout << "NONE" << " ";
|
||||
|
||||
std::cout <<"\n"<<std::endl;
|
||||
}
|
||||
|
||||
std::cout << "\nNote: For non-supported peer2peer devices, memcopy will use/follow the normal behaviour (GPU1-->host then host-->GPU2)\n\n" << std::endl;
|
||||
}
|
||||
|
||||
void enablePeer2Peer(int currentGpu, int peerGpu)
|
||||
{
|
||||
int canAccessPeer;
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
|
||||
|
||||
if(canAccessPeer==1){
|
||||
hipDeviceEnablePeerAccess(peerGpu, 0);
|
||||
}
|
||||
}
|
||||
|
||||
void disablePeer2Peer(int currentGpu, int peerGpu)
|
||||
{
|
||||
int canAccessPeer;
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
|
||||
|
||||
if(canAccessPeer==1){
|
||||
hipDeviceDisablePeerAccess(peerGpu);
|
||||
}
|
||||
}
|
||||
|
||||
std::string gpuIDToString(int gpuID)
|
||||
{
|
||||
using namespace std;
|
||||
stringstream ss;
|
||||
ss << gpuID;
|
||||
return ss.str();
|
||||
}
|
||||
|
||||
void RunBenchmark_P2P_Unidir(ResultDatabase &resultDB)
|
||||
{
|
||||
int gpuCount;
|
||||
hipGetDeviceCount(&gpuCount);
|
||||
|
||||
int currentGpu, peerGpu;
|
||||
|
||||
long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
|
||||
|
||||
for (currentGpu=0; currentGpu<gpuCount; currentGpu++) {
|
||||
|
||||
for (peerGpu=0; peerGpu<gpuCount; peerGpu++){
|
||||
|
||||
if (currentGpu == peerGpu)
|
||||
continue;
|
||||
|
||||
float *currentGpuMem, *peerGpuMem;
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipMalloc((void**)¤tGpuMem, sizeof(float) * numMaxFloats);
|
||||
|
||||
hipSetDevice(peerGpu);
|
||||
hipMalloc((void**)&peerGpuMem, sizeof(float) * numMaxFloats);
|
||||
|
||||
enablePeer2Peer(currentGpu, peerGpu);
|
||||
|
||||
hipEvent_t start, stop;
|
||||
hipEventCreate(&start);
|
||||
hipEventCreate(&stop);
|
||||
CHECK_HIP_ERROR();
|
||||
|
||||
// Three passes, forward and backward both
|
||||
for (int pass = 0; pass < p_iterations; pass++)
|
||||
{
|
||||
// store the times temporarily to estimate latency
|
||||
//float times[nSizes];
|
||||
// Step through sizes forward on even passes and backward on odd
|
||||
for (int i = 0; i < nSizes; i++)
|
||||
{
|
||||
int sizeIndex;
|
||||
if ((pass % 2) == 0)
|
||||
sizeIndex = i;
|
||||
else
|
||||
sizeIndex = (nSizes - 1) - i;
|
||||
|
||||
const int thisSize = p_onesize ? p_onesize : sizes[sizeIndex];
|
||||
const int nbytes = sizeToBytes(thisSize);
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
hipEventRecord(start, 0);
|
||||
|
||||
for (int j=0;j<p_beatsperiteration;j++) {
|
||||
hipMemcpy(peerGpuMem, currentGpuMem, nbytes, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
hipEventRecord(stop, 0);
|
||||
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
float t = 0;
|
||||
hipEventElapsedTime(&t, start, stop);
|
||||
//times[sizeIndex] = t;
|
||||
|
||||
// Convert to GB/sec
|
||||
if (p_verbose)
|
||||
{
|
||||
std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n";
|
||||
}
|
||||
|
||||
double speed = (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000*1000)) / t;
|
||||
char sizeStr[256];
|
||||
if (p_beatsperiteration>1) {
|
||||
sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration);
|
||||
} else {
|
||||
sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str());
|
||||
}
|
||||
|
||||
string cGpu, pGpu;
|
||||
cGpu = gpuIDToString(currentGpu);
|
||||
pGpu = gpuIDToString(peerGpu);
|
||||
|
||||
resultDB.AddResult(std::string("p2p_uni") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "GB/sec", speed);
|
||||
resultDB.AddResult(std::string("P2P_uni") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "ms", t);
|
||||
|
||||
if (p_onesize) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
if (p_onesize) {
|
||||
numMaxFloats = sizeToBytes(p_onesize) / sizeof(float);
|
||||
}
|
||||
|
||||
disablePeer2Peer(currentGpu, peerGpu);
|
||||
|
||||
hipEventDestroy(start);
|
||||
hipEventDestroy(stop);
|
||||
|
||||
// Cleanup
|
||||
hipFree((void*)currentGpuMem);
|
||||
hipFree((void*)peerGpuMem);
|
||||
CHECK_HIP_ERROR();
|
||||
|
||||
hipSetDevice(peerGpu);
|
||||
hipDeviceReset();
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipDeviceReset();
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) {
|
||||
|
||||
int gpuCount;
|
||||
hipGetDeviceCount(&gpuCount);
|
||||
|
||||
hipStream_t stream[2];
|
||||
|
||||
int currentGpu, peerGpu;
|
||||
|
||||
long long numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
|
||||
|
||||
for (currentGpu=0; currentGpu<gpuCount; currentGpu++) {
|
||||
|
||||
for (peerGpu=0; peerGpu<gpuCount; peerGpu++){
|
||||
|
||||
if (currentGpu == peerGpu)
|
||||
continue;
|
||||
|
||||
float *currentGpuMem[2], *peerGpuMem[2];
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipMalloc((void**)¤tGpuMem[0], sizeof(float) * numMaxFloats);
|
||||
hipMalloc((void**)¤tGpuMem[1], sizeof(float) * numMaxFloats);
|
||||
|
||||
hipSetDevice(peerGpu);
|
||||
hipMalloc((void**)&peerGpuMem[0], sizeof(float) * numMaxFloats);
|
||||
hipMalloc((void**)&peerGpuMem[1], sizeof(float) * numMaxFloats);
|
||||
|
||||
enablePeer2Peer(currentGpu, peerGpu);
|
||||
|
||||
hipEvent_t start, stop;
|
||||
hipEventCreate(&start);
|
||||
hipEventCreate(&stop);
|
||||
CHECK_HIP_ERROR();
|
||||
|
||||
hipStreamCreate(&stream[0]);
|
||||
hipStreamCreate(&stream[1]);
|
||||
|
||||
// Three passes, forward and backward both
|
||||
for (int pass = 0; pass < p_iterations; pass++)
|
||||
{
|
||||
// store the times temporarily to estimate latency
|
||||
//float times[nSizes];
|
||||
// Step through sizes forward on even passes and backward on odd
|
||||
for (int i = 0; i < nSizes; i++)
|
||||
{
|
||||
int sizeIndex;
|
||||
if ((pass % 2) == 0)
|
||||
sizeIndex = i;
|
||||
else
|
||||
sizeIndex = (nSizes - 1) - i;
|
||||
|
||||
const int thisSize = p_onesize ? p_onesize : sizes[sizeIndex];
|
||||
const int nbytes = sizeToBytes(thisSize);
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
hipEventRecord(start, 0);
|
||||
|
||||
for (int j=0;j<p_beatsperiteration;j++) {
|
||||
hipMemcpyAsync(peerGpuMem[0], currentGpuMem[0], nbytes, hipMemcpyDeviceToDevice, stream[0]);
|
||||
hipMemcpyAsync(currentGpuMem[1], peerGpuMem[1], nbytes, hipMemcpyDeviceToDevice, stream[1]);
|
||||
}
|
||||
|
||||
hipEventRecord(stop, 0);
|
||||
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
float t = 0;
|
||||
hipEventElapsedTime(&t, start, stop);
|
||||
//times[sizeIndex] = t;
|
||||
|
||||
// Convert to GB/sec
|
||||
if (p_verbose)
|
||||
{
|
||||
std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n";
|
||||
}
|
||||
|
||||
double speed = (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000*1000)) / t;
|
||||
char sizeStr[256];
|
||||
if (p_beatsperiteration>1) {
|
||||
sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration);
|
||||
} else {
|
||||
sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str());
|
||||
}
|
||||
|
||||
string cGpu, pGpu;
|
||||
cGpu = gpuIDToString(currentGpu);
|
||||
pGpu = gpuIDToString(peerGpu);
|
||||
|
||||
resultDB.AddResult(std::string("p2p_bi") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "GB/sec", speed);
|
||||
resultDB.AddResult(std::string("P2P_bi") + "_gpu" + std::string(cGpu)+ "_gpu" + std::string(pGpu), sizeStr, "ms", t);
|
||||
|
||||
if (p_onesize) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
if (p_onesize) {
|
||||
numMaxFloats = sizeToBytes(p_onesize) / sizeof(float);
|
||||
}
|
||||
|
||||
disablePeer2Peer(currentGpu, peerGpu);
|
||||
|
||||
hipEventDestroy(start);
|
||||
hipEventDestroy(stop);
|
||||
|
||||
for (int i=0; i<2; i++) {
|
||||
hipStreamDestroy(stream[i]);
|
||||
|
||||
hipFree((void*)currentGpuMem[i]);
|
||||
hipFree((void*)peerGpuMem[i]);
|
||||
CHECK_HIP_ERROR();
|
||||
}
|
||||
|
||||
hipSetDevice(peerGpu);
|
||||
hipDeviceReset();
|
||||
|
||||
hipSetDevice(currentGpu);
|
||||
hipDeviceReset();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void printConfig() {
|
||||
hipDeviceProp_t props;
|
||||
hipGetDeviceProperties(&props, p_device);
|
||||
@@ -662,9 +987,9 @@ void help() {
|
||||
printf (" --d2h : Run only device-to-host test.\n");
|
||||
printf (" --h2d : Run only host-to-device test.\n");
|
||||
printf (" --bidir : Run only bidir copy test.\n");
|
||||
printf (" --p2p : Run only peer2peer unidir and bidir copy tests.\n");
|
||||
printf (" --verbose : Print verbose status messages as test is run.\n");
|
||||
printf (" --detailed : Print detailed report (including all trials).\n");
|
||||
|
||||
printf (" --async : Use hipMemcpyAsync(with NULL stream) for H2D/D2H. Default uses hipMemcpy.\n");
|
||||
printf (" --onesize, -o : Only run one measurement, at specified size (in KB, or if negative in bytes)\n");
|
||||
|
||||
@@ -712,6 +1037,12 @@ int parseStandardArguments(int argc, char *argv[])
|
||||
p_d2h = false;
|
||||
p_bidir = true;
|
||||
|
||||
} else if (!strcmp(arg, "--p2p")) {
|
||||
p_h2d = false;
|
||||
p_d2h = false;
|
||||
p_bidir = false;
|
||||
p_p2p = true;
|
||||
|
||||
} else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) {
|
||||
help();
|
||||
exit(EXIT_SUCCESS);
|
||||
@@ -737,39 +1068,57 @@ int main(int argc, char *argv[])
|
||||
{
|
||||
parseStandardArguments(argc, argv);
|
||||
|
||||
printConfig();
|
||||
if (p_p2p) {
|
||||
checkPeer2PeerSupport();
|
||||
|
||||
if (p_h2d) {
|
||||
ResultDatabase resultDB;
|
||||
RunBenchmark_H2D(resultDB);
|
||||
ResultDatabase resultDB_Unidir, resultDB_Bidir;
|
||||
|
||||
resultDB.DumpSummary(std::cout);
|
||||
RunBenchmark_P2P_Unidir(resultDB_Unidir);
|
||||
RunBenchmark_P2P_Bidir(resultDB_Bidir);
|
||||
|
||||
resultDB_Unidir.DumpSummary(std::cout);
|
||||
resultDB_Bidir.DumpSummary(std::cout);
|
||||
|
||||
if (p_detailed) {
|
||||
resultDB.DumpDetailed(std::cout);
|
||||
resultDB_Unidir.DumpDetailed(std::cout);
|
||||
resultDB_Bidir.DumpDetailed(std::cout);
|
||||
}
|
||||
}
|
||||
else {
|
||||
printConfig();
|
||||
|
||||
if (p_d2h) {
|
||||
ResultDatabase resultDB;
|
||||
RunBenchmark_D2H(resultDB);
|
||||
if (p_h2d) {
|
||||
ResultDatabase resultDB;
|
||||
RunBenchmark_H2D(resultDB);
|
||||
|
||||
resultDB.DumpSummary(std::cout);
|
||||
resultDB.DumpSummary(std::cout);
|
||||
|
||||
if (p_detailed) {
|
||||
resultDB.DumpDetailed(std::cout);
|
||||
if (p_detailed) {
|
||||
resultDB.DumpDetailed(std::cout);
|
||||
}
|
||||
}
|
||||
|
||||
if (p_d2h) {
|
||||
ResultDatabase resultDB;
|
||||
RunBenchmark_D2H(resultDB);
|
||||
|
||||
resultDB.DumpSummary(std::cout);
|
||||
|
||||
if (p_detailed) {
|
||||
resultDB.DumpDetailed(std::cout);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
if (p_bidir) {
|
||||
ResultDatabase resultDB;
|
||||
RunBenchmark_Bidir(resultDB);
|
||||
if (p_bidir) {
|
||||
ResultDatabase resultDB;
|
||||
RunBenchmark_Bidir(resultDB);
|
||||
|
||||
resultDB.DumpSummary(std::cout);
|
||||
resultDB.DumpSummary(std::cout);
|
||||
|
||||
if (p_detailed) {
|
||||
resultDB.DumpDetailed(std::cout);
|
||||
if (p_detailed) {
|
||||
resultDB.DumpDetailed(std::cout);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,35 @@
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
TARGET=hcc
|
||||
|
||||
SOURCES = inline_asm.cpp
|
||||
OBJECTS = $(SOURCES:.cpp=.o)
|
||||
|
||||
EXECUTABLE=./inline_asm
|
||||
|
||||
.PHONY: test
|
||||
|
||||
|
||||
all: $(EXECUTABLE) test
|
||||
|
||||
CXXFLAGS =-g
|
||||
CXX=$(HIPCC)
|
||||
|
||||
|
||||
$(EXECUTABLE): $(OBJECTS)
|
||||
$(HIPCC) $(OBJECTS) -o $@
|
||||
|
||||
|
||||
test: $(EXECUTABLE)
|
||||
$(EXECUTABLE)
|
||||
|
||||
|
||||
clean:
|
||||
rm -f $(EXECUTABLE)
|
||||
rm -f $(OBJECTS)
|
||||
|
||||
@@ -0,0 +1,47 @@
|
||||
## inline asm ###
|
||||
|
||||
This tutorial is about how to use inline GCN asm in kernel. In this tutorial, we'll explain how to by using the simple Matrix Transpose.
|
||||
|
||||
## Introduction:
|
||||
|
||||
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
|
||||
[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/)
|
||||
|
||||
For more information:
|
||||
[AMD GCN3 ISA Architecture Manual](http://gpuopen.com/compute-product/amd-gcn3-isa-architecture-manual/)
|
||||
[User Guide for AMDGPU Back-end](llvm.org/docs/AMDGPUUsage.html)
|
||||
|
||||
## Requirement:
|
||||
For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-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
|
||||
|
||||
We will be using the Simple Matrix Transpose application from the our very first tutorial.
|
||||
|
||||
## asm() Assembler statement
|
||||
|
||||
We insert the GCN isa into the kernel using asm() Assembler statement. In the same sourcecode, we used for MatrixTranspose. We'll add the following:
|
||||
|
||||
` asm volatile ("v_mov_b32_e32 %0, %1" : "=v" (out[x*width + y]) : "v" (in[y*width + x])); `
|
||||
|
||||
## How to build and run:
|
||||
Use the make command and execute it using ./exe
|
||||
Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.
|
||||
|
||||
|
||||
## More Info:
|
||||
- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md)
|
||||
- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md)
|
||||
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
|
||||
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
|
||||
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
|
||||
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
|
||||
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
|
||||
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
|
||||
@@ -0,0 +1,174 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include<iostream>
|
||||
|
||||
// hip header file
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define WIDTH 1024
|
||||
|
||||
#define NUM (WIDTH*WIDTH)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 4
|
||||
#define THREADS_PER_BLOCK_Y 4
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
// Device (Kernel) function, it must be void
|
||||
// hipLaunchParm provides the execution configuration
|
||||
__global__ void matrixTranspose(hipLaunchParm lp,
|
||||
float *out,
|
||||
float *in,
|
||||
const int width)
|
||||
{
|
||||
|
||||
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
asm volatile ("v_mov_b32_e32 %0, %1" : "=v" (out[x*width + y]) : "v" (in[y*width + x]));
|
||||
}
|
||||
|
||||
// CPU implementation of matrix transpose
|
||||
void matrixTransposeCPUReference(
|
||||
float * output,
|
||||
float * input,
|
||||
const unsigned int width)
|
||||
{
|
||||
for(unsigned int j=0; j < width; j++)
|
||||
{
|
||||
for(unsigned int i=0; i < width; i++)
|
||||
{
|
||||
output[i*width + j] = input[j*width + i];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
|
||||
float* Matrix;
|
||||
float* TransposeMatrix;
|
||||
float* cpuTransposeMatrix;
|
||||
|
||||
float* gpuMatrix;
|
||||
float* gpuTransposeMatrix;
|
||||
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
|
||||
std::cout << "Device name " << devProp.name << std::endl;
|
||||
|
||||
hipEvent_t start, stop;
|
||||
hipEventCreate(&start);
|
||||
hipEventCreate(&stop);
|
||||
float eventMs = 1.0f;
|
||||
|
||||
int i;
|
||||
int errors;
|
||||
|
||||
Matrix = (float*)malloc(NUM * sizeof(float));
|
||||
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
|
||||
// initialize the input data
|
||||
for (i = 0; i < NUM; i++) {
|
||||
Matrix[i] = (float)i*10.0f;
|
||||
}
|
||||
|
||||
// allocate the memory on the device side
|
||||
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
|
||||
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
|
||||
|
||||
// Record the start event
|
||||
hipEventRecord(start, NULL);
|
||||
|
||||
// Memory transfer from host to device
|
||||
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
|
||||
|
||||
// Record the stop event
|
||||
hipEventRecord(stop, NULL);
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
hipEventElapsedTime(&eventMs, start, stop);
|
||||
|
||||
printf ("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs);
|
||||
|
||||
// Record the start event
|
||||
hipEventRecord(start, NULL);
|
||||
|
||||
// Lauching kernel from host
|
||||
hipLaunchKernel(matrixTranspose,
|
||||
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
|
||||
0, 0,
|
||||
gpuTransposeMatrix , gpuMatrix, WIDTH);
|
||||
|
||||
// Record the stop event
|
||||
hipEventRecord(stop, NULL);
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
hipEventElapsedTime(&eventMs, start, stop);
|
||||
|
||||
printf ("kernel Execution time = %6.3fms\n", eventMs);
|
||||
|
||||
// Record the start event
|
||||
hipEventRecord(start, NULL);
|
||||
|
||||
// Memory transfer from device to host
|
||||
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
|
||||
|
||||
// Record the stop event
|
||||
hipEventRecord(stop, NULL);
|
||||
hipEventSynchronize(stop);
|
||||
|
||||
hipEventElapsedTime(&eventMs, start, stop);
|
||||
|
||||
printf ("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs);
|
||||
|
||||
// CPU MatrixTranspose computation
|
||||
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
|
||||
|
||||
// verify the results
|
||||
errors = 0;
|
||||
double eps = 1.0E-6;
|
||||
for (i = 0; i < NUM; i++) {
|
||||
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
|
||||
printf("gpu%f cpu %f \n",TransposeMatrix[i],cpuTransposeMatrix[i]);
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors!=0) {
|
||||
printf("FAILED: %d errors\n",errors);
|
||||
} else {
|
||||
printf ("PASSED!\n");
|
||||
}
|
||||
|
||||
//free the resources on device side
|
||||
hipFree(gpuMatrix);
|
||||
hipFree(gpuTransposeMatrix);
|
||||
|
||||
//free the resources on host side
|
||||
free(Matrix);
|
||||
free(TransposeMatrix);
|
||||
free(cpuTransposeMatrix);
|
||||
|
||||
return errors;
|
||||
}
|
||||
@@ -55,13 +55,9 @@ void checkPeer2PeerSupport()
|
||||
{
|
||||
int gpuCount;
|
||||
int canAccessPeer;
|
||||
int p2pCapableDeviceCount=0;
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&gpuCount));
|
||||
|
||||
if (gpuCount < 2)
|
||||
printf("Peer2Peer application requires atleast 2 gpu devices");
|
||||
|
||||
for (int currentGpu=0; currentGpu<gpuCount; currentGpu++)
|
||||
{
|
||||
HIPCHECK(hipSetDevice(currentGpu));
|
||||
@@ -161,6 +157,12 @@ int main(){
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&gpuCount));
|
||||
|
||||
if (gpuCount < 2)
|
||||
{
|
||||
printf("Peer2Peer application requires atleast 2 gpu devices");
|
||||
return 0;
|
||||
}
|
||||
|
||||
currentGpu = 0;
|
||||
peerGpu = (currentGpu + 1);
|
||||
|
||||
|
||||
@@ -0,0 +1,39 @@
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
ifeq (gfx701, $(findstring gfx701,$(HCC_AMDGPU_TARGET)))
|
||||
$(error gfx701 is not a supported device for this sample)
|
||||
endif
|
||||
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
TARGET=hcc
|
||||
|
||||
SOURCES = unroll.cpp
|
||||
OBJECTS = $(SOURCES:.cpp=.o)
|
||||
|
||||
EXECUTABLE=./unroll
|
||||
|
||||
.PHONY: test
|
||||
|
||||
|
||||
all: $(EXECUTABLE) test
|
||||
|
||||
CXXFLAGS =-g
|
||||
CXX=$(HIPCC)
|
||||
|
||||
|
||||
$(EXECUTABLE): $(OBJECTS)
|
||||
$(HIPCC) $(OBJECTS) -o $@
|
||||
|
||||
|
||||
test: $(EXECUTABLE)
|
||||
$(EXECUTABLE)
|
||||
|
||||
|
||||
clean:
|
||||
rm -f $(EXECUTABLE)
|
||||
rm -f $(OBJECTS)
|
||||
rm -f $(HIP_PATH)/src/*.o
|
||||
@@ -0,0 +1,48 @@
|
||||
## Using Pragma unroll ###
|
||||
|
||||
In this tutorial, we'll explain how to use #pragma unroll to improve the performance.
|
||||
|
||||
## Introduction:
|
||||
|
||||
Loop unrolling optimization hints can be specified with #pragma unroll and #pragma nounroll. The pragma is placed immediately before a for loop.
|
||||
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/GPUOpen-ProfessionalCompute-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
|
||||
|
||||
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); `
|
||||
` } `
|
||||
|
||||
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.
|
||||
|
||||
## How to build and run:
|
||||
Use the make command and execute it using ./exe
|
||||
Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.
|
||||
|
||||
## requirement for nvidia
|
||||
please make sure you have a 3.0 or higher compute capable device in order to use warp shfl operations and add `-gencode arch=compute=30, code=sm_30` nvcc flag in the Makefile while using this application.
|
||||
|
||||
## More Info:
|
||||
- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md)
|
||||
- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md)
|
||||
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
|
||||
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
|
||||
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
|
||||
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
|
||||
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
|
||||
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
|
||||
@@ -0,0 +1,141 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include<iostream>
|
||||
|
||||
// hip header file
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
|
||||
#define WIDTH 4
|
||||
|
||||
#define NUM (WIDTH*WIDTH)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 4
|
||||
#define THREADS_PER_BLOCK_Y 4
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
// Device (Kernel) function, it must be void
|
||||
// hipLaunchParm provides the execution configuration
|
||||
__global__ void matrixTranspose(hipLaunchParm lp,
|
||||
float *out,
|
||||
float *in,
|
||||
const int width)
|
||||
{
|
||||
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
float val = in[x];
|
||||
|
||||
#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);
|
||||
}
|
||||
}
|
||||
|
||||
// CPU implementation of matrix transpose
|
||||
void matrixTransposeCPUReference(
|
||||
float * output,
|
||||
float * input,
|
||||
const unsigned int width)
|
||||
{
|
||||
for(unsigned int j=0; j < width; j++)
|
||||
{
|
||||
for(unsigned int i=0; i < width; i++)
|
||||
{
|
||||
output[i*width + j] = input[j*width + i];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
|
||||
float* Matrix;
|
||||
float* TransposeMatrix;
|
||||
float* cpuTransposeMatrix;
|
||||
|
||||
float* gpuMatrix;
|
||||
float* gpuTransposeMatrix;
|
||||
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
|
||||
std::cout << "Device name " << devProp.name << std::endl;
|
||||
|
||||
int i;
|
||||
int errors;
|
||||
|
||||
Matrix = (float*)malloc(NUM * sizeof(float));
|
||||
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
|
||||
// initialize the input data
|
||||
for (i = 0; i < NUM; i++) {
|
||||
Matrix[i] = (float)i*10.0f;
|
||||
}
|
||||
|
||||
// allocate the memory on the device side
|
||||
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
|
||||
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
|
||||
|
||||
// Memory transfer from host to device
|
||||
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
|
||||
|
||||
// Lauching kernel from host
|
||||
hipLaunchKernel(matrixTranspose,
|
||||
dim3(1),
|
||||
dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y),
|
||||
0, 0,
|
||||
gpuTransposeMatrix , gpuMatrix, WIDTH);
|
||||
|
||||
// Memory transfer from device to host
|
||||
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
|
||||
|
||||
// CPU MatrixTranspose computation
|
||||
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
|
||||
|
||||
// verify the results
|
||||
errors = 0;
|
||||
double eps = 1.0E-6;
|
||||
for (i = 0; i < NUM; i++) {
|
||||
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
|
||||
printf("%d cpu: %f gpu %f\n",i,cpuTransposeMatrix[i],TransposeMatrix[i]);
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors!=0) {
|
||||
printf("FAILED: %d errors\n",errors);
|
||||
} else {
|
||||
printf ("PASSED!\n");
|
||||
}
|
||||
|
||||
//free the resources on device side
|
||||
hipFree(gpuMatrix);
|
||||
hipFree(gpuTransposeMatrix);
|
||||
|
||||
//free the resources on host side
|
||||
free(Matrix);
|
||||
free(TransposeMatrix);
|
||||
free(cpuTransposeMatrix);
|
||||
|
||||
return errors;
|
||||
}
|
||||
Reference in New Issue
Block a user