AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

Apply .clangformat to all repo source files

Procházet zdrojové kódy 
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344


[ROCm/hip-tests commit: 6b09bde675]
Tento commit je obsažen v:
Maneesh Gupta
2018-03-12 11:29:03 +05:30
rodič dd60eaa156
revize 3277453ea5
39 změnil soubory, kde provedl 2386 přidání a 2967 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp
+32 -35
Zobrazit soubor
@@ -28,79 +28,76 @@ THE SOFTWARE.
#endif
#define CHECK(cmd) \
{\
hipError_t error = cmd;\
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
exit(EXIT_FAILURE);\
}\
}
#define CHECK(cmd) \
{ \
hipError_t error = cmd; \
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \
__FILE__, __LINE__); \
exit(EXIT_FAILURE); \
} \
}
__global__ void
bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N)
{
__global__ void bit_extract_kernel(hipLaunchParm lp, uint32_t* C_d, const uint32_t* A_d, size_t N) {
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x ;
size_t stride = hipBlockDim_x * hipGridDim_x;
for (size_t i=offset; i<N; i+=stride) {
for (size_t i = offset; i < N; i += stride) {
#ifdef __HIP_PLATFORM_HCC__
C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4);
#else /* defined __HIP_PLATFORM_NVCC__ or other path */
C_d[i] = ((A_d[i] & 0xf00) >> 8);
C_d[i] = ((A_d[i] & 0xf00) >> 8);
#endif
}
}
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
uint32_t *A_d, *C_d;
uint32_t *A_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(uint32_t);
int deviceId;
CHECK (hipGetDevice(&deviceId));
CHECK(hipGetDevice(&deviceId));
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, deviceId));
printf ("info: running on device #%d %s\n", deviceId, props.name);
printf("info: running on device #%d %s\n", deviceId, props.name);
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
A_h = (uint32_t*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
C_h = (uint32_t*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
for (size_t i=0; i<N; i++)
{
for (size_t i = 0; i < N; i++) {
A_h[i] = i;
}
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
printf ("info: copy Host2Device\n");
CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
printf("info: copy Host2Device\n");
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
printf ("info: launch 'bit_extract_kernel' \n");
printf("info: launch 'bit_extract_kernel' \n");
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: copy Device2Host\n");
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf("info: copy Device2Host\n");
CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf ("info: check result\n");
for (size_t i=0; i<N; i++) {
printf("info: check result\n");
for (size_t i = 0; i < N; i++) {
unsigned Agold = ((A_h[i] & 0xf00) >> 8);
if (C_h[i] != Agold) {
fprintf (stderr, "mismatch detected.\n");
printf ("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]);
fprintf(stderr, "mismatch detected.\n");
printf("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]);
CHECK(hipErrorUnknown);
}
}
printf ("PASSED!\n");
printf("PASSED!\n");
}
projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp
+19 -18
Zobrazit soubor
@@ -26,8 +26,8 @@ THE SOFTWARE.
// will automatically copy data to and from the host, without the user needing
// to manually perform such copies. This is an excellent mode for developers
// new to GPU programming and matches the memory models provided by recent systems where
// CPU and GPU share the same memory pool. Advanced programmers may prefer
// more explicit control over the data movement - shown in the other vadd_hc_array and
// CPU and GPU share the same memory pool. Advanced programmers may prefer
// more explicit control over the data movement - shown in the other vadd_hc_array and
// vadd_hc_am examples.
// This example shows the similarity between C++AMP and and HC for simple cases where
// implicit data transfer is used - really the only difference is the namespace.
@@ -35,8 +35,7 @@ THE SOFTWARE.
#include <amp.h>
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
int sizeElements = 1000000;
bool pass = true;
@@ -46,28 +45,30 @@ int main(int argc, char *argv[])
concurrency::array_view<float> C(sizeElements);
// Initialize host data
for (int i=0; i<sizeElements; i++) {
A[i] = 1.618f * i;
for (int i = 0; i < sizeElements; i++) {
A[i] = 1.618f * i;
B[i] = 3.142f * i;
}
C.discard_data(); // tell runtime not to copy CPU host data.
C.discard_data(); // tell runtime not to copy CPU host data.
// Launch kernel onto default accelerator
// The HCC runtime will ensure that A and B are available on the accelerator before launching the kernel.
concurrency::parallel_for_each(concurrency::extent<1> (sizeElements),
[=] (concurrency::index<1> idx) restrict(amp) {
int i = idx[0];
C[i] = A[i] + B[i];
});
// The HCC runtime will ensure that A and B are available on the accelerator before launching
// the kernel.
concurrency::parallel_for_each(concurrency::extent<1>(sizeElements),
[=](concurrency::index<1> idx) restrict(amp) {
int i = idx[0];
C[i] = A[i] + B[i];
});
for (int i=0; i<sizeElements; i++) {
float ref= 1.618f * i + 3.142f * i;
// Because C is an array_view, the HCC runtime will copy C back to host at first access here:
for (int i = 0; i < sizeElements; i++) {
float ref = 1.618f * i + 3.142f * i;
// Because C is an array_view, the HCC runtime will copy C back to host at first access
// here:
if (C[i] != ref) {
printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
pass = false;
}
};
if (pass) printf ("PASSED!\n");
if (pass) printf("PASSED!\n");
}
projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
+26 -26
Zobrazit soubor
@@ -24,21 +24,20 @@ THE SOFTWARE.
// AM provides a set of c-style memory management routines for allocating,
// freeing, and copying memory. am_alloc returns a device pointer
// which can only be used on the device. The programmer has full control
// over when data is copied.
// over when data is copied.
#include <hc.hpp>
#include <hc_am.hpp>
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
int sizeElements = 1000000;
size_t sizeBytes = sizeElements * sizeof(float);
bool pass = true;
// Allocate host memory
float *A_h = (float*)malloc(sizeBytes);
float *B_h = (float*)malloc(sizeBytes);
float *C_h = (float*)malloc(sizeBytes);
float* A_h = (float*)malloc(sizeBytes);
float* B_h = (float*)malloc(sizeBytes);
float* C_h = (float*)malloc(sizeBytes);
// Allocate device pointers:
// Unlike array_view, these must be explicitly managed by user:
@@ -51,36 +50,37 @@ int main(int argc, char *argv[])
C_d = hc::am_alloc(sizeBytes, acc, 0);
// Initialize host data
for (int i=0; i<sizeElements; i++) {
A_h[i] = 1.618f * i;
for (int i = 0; i < sizeElements; i++) {
A_h[i] = 1.618f * i;
B_h[i] = 3.142f * i;
C_h[i] = 0;
}
av.copy(A_h, A_d, sizeBytes); // C++ copy H2D
av.copy(B_h, B_d, sizeBytes); // C++ copy H2D
av.copy(A_h, A_d, sizeBytes); // C++ copy H2D
av.copy(B_h, B_d, sizeBytes); // C++ copy H2D
// Launch kernel onto AV.
// Launch kernel onto AV.
// Because the kernel PFE and the copies are submitted to same AV, they will execute in order
// and we don't need additional synchronization to ensure the copies complete before the PFE begins.
hc::completion_future cf=
hc::parallel_for_each(av, hc::extent<1> (sizeElements),
[=] (hc::index<1> idx) [[hc]] {
int i = idx[0];
C_d[i] = A_d[i] + B_d[i];
});
// This copy is in same AV as the kernel and thus will wait for the kernel to finish before executing.
av.copy(C_d, C_h, sizeBytes); // C++ copy D2H
// and we don't need additional synchronization to ensure the copies complete before the PFE
// begins.
hc::completion_future cf =
hc::parallel_for_each(av, hc::extent<1>(sizeElements), [=](hc::index<1> idx)[[hc]] {
int i = idx[0];
C_d[i] = A_d[i] + B_d[i];
});
for (int i=0; i<sizeElements; i++) {
float ref= 1.618f * i + 3.142f * i;
// This copy is in same AV as the kernel and thus will wait for the kernel to finish before
// executing.
av.copy(C_d, C_h, sizeBytes); // C++ copy D2H
for (int i = 0; i < sizeElements; i++) {
float ref = 1.618f * i + 3.142f * i;
if (C_h[i] != ref) {
printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
pass = false;
}
};
if (pass) printf ("PASSED!\n");
if (pass) printf("PASSED!\n");
}
projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
+17 -18
Zobrazit soubor
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
// Simple test showing how to use HC syntax with array.
// Array provides a type-safe C++ mechanism to allocate accelerator memory.
// Array provides a type-safe C++ mechanism to allocate accelerator memory.
// Like array_view, hc::array provides multi-dimensional indexing capability,
// and is typed. However, unlike array_view, hc::array does not provide
// automatic data management capabilities - instead the programmer
@@ -29,16 +29,15 @@ THE SOFTWARE.
#include <hc.hpp>
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
int sizeElements = 1000000;
size_t sizeBytes = sizeElements * sizeof(float);
bool pass = true;
// Allocate host memory
float *A_h = (float*)malloc(sizeBytes);
float *B_h = (float*)malloc(sizeBytes);
float *C_h = (float*)malloc(sizeBytes);
float* A_h = (float*)malloc(sizeBytes);
float* B_h = (float*)malloc(sizeBytes);
float* C_h = (float*)malloc(sizeBytes);
// Allocate device arrays<>
// Unlike array_view, these must be explicitly managed by user:
@@ -47,32 +46,32 @@ int main(int argc, char *argv[])
hc::array<float> C_d(sizeElements);
// Initialize host data
for (int i=0; i<sizeElements; i++) {
A_h[i] = 1.618f * i;
for (int i = 0; i < sizeElements; i++) {
A_h[i] = 1.618f * i;
B_h[i] = 3.142f * i;
}
hc::copy(A_h, A_d); // C++ copy H2D
hc::copy(B_h, B_d); // C++ copy H2D
hc::copy(A_h, A_d); // C++ copy H2D
hc::copy(B_h, B_d); // C++ copy H2D
// Launch kernel onto default accelerator:
// array<> types are not implicitly copied, so we performed copies above.
hc::parallel_for_each(hc::extent<1> (sizeElements),
[&] (hc::index<1> idx) [[hc]] {
hc::parallel_for_each(hc::extent<1>(sizeElements), [&](hc::index<1> idx)[[hc]] {
int i = idx[0];
C_d[i] = A_d[i] + B_d[i];
});
// HCC runtime knows that C_d depends on previous PFE and will force the copy to wait for the PFE to complte.
hc::copy(C_d, C_h); // C++ copy D2H
// HCC runtime knows that C_d depends on previous PFE and will force the copy to wait for the
// PFE to complte.
hc::copy(C_d, C_h); // C++ copy D2H
for (int i=0; i<sizeElements; i++) {
float ref= 1.618f * i + 3.142f * i;
for (int i = 0; i < sizeElements; i++) {
float ref = 1.618f * i + 3.142f * i;
if (C_h[i] != ref) {
printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
pass = false;
}
};
if (pass) printf ("PASSED!\n");
if (pass) printf("PASSED!\n");
}
projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
+15 -15
Zobrazit soubor
@@ -26,8 +26,8 @@ THE SOFTWARE.
// will automatically copy data to and from the host, without the user needing
// to manually perform such copies. This is an excellent mode for developers
// new to GPU programming and matches the memory models provided by recent systems where
// CPU and GPU share the same memory pool. Advanced programmers may prefer
// more explicit control over the data movement - shown in the other vadd_hc_array and
// CPU and GPU share the same memory pool. Advanced programmers may prefer
// more explicit control over the data movement - shown in the other vadd_hc_array and
// vadd_hc_am examples.
// This example shows the similarity between C++AMP and and HC for simple cases where
// implicit data transfer is used - really the only difference is the namespace.
@@ -35,8 +35,7 @@ THE SOFTWARE.
#include <hc.hpp>
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
int sizeElements = 1000000;
bool pass = true;
@@ -46,28 +45,29 @@ int main(int argc, char *argv[])
hc::array_view<float> C(sizeElements);
// Initialize host data
for (int i=0; i<sizeElements; i++) {
A[i] = 1.618f * i;
for (int i = 0; i < sizeElements; i++) {
A[i] = 1.618f * i;
B[i] = 3.142f * i;
}
C.discard_data(); // tell runtime not to copy CPU host data.
C.discard_data(); // tell runtime not to copy CPU host data.
// Launch kernel onto default accelerator:
// The HCC runtime will ensure that A and B are available on the accelerator before launching the kernel.
hc::parallel_for_each(hc::extent<1> (sizeElements),
[=] (hc::index<1> idx) [[hc]] {
// The HCC runtime will ensure that A and B are available on the accelerator before launching
// the kernel.
hc::parallel_for_each(hc::extent<1>(sizeElements), [=](hc::index<1> idx)[[hc]] {
int i = idx[0];
C[i] = A[i] + B[i];
});
for (int i=0; i<sizeElements; i++) {
float ref= 1.618f * i + 3.142f * i;
// Because C is an array_view, the HCC runtime will copy C back to host at first access here:
for (int i = 0; i < sizeElements; i++) {
float ref = 1.618f * i + 3.142f * i;
// Because C is an array_view, the HCC runtime will copy C back to host at first access
// here:
if (C[i] != ref) {
printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
pass = false;
}
};
if (pass) printf ("PASSED!\n");
if (pass) printf("PASSED!\n");
}
projects/hip-tests/samples/0_Intro/hcc_dialects/vadd_hip.cpp
+13 -15
Zobrazit soubor
@@ -22,8 +22,7 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
__global__ void vadd_hip(hipLaunchParm lp, const float *a, const float *b, float *c, int N)
{
__global__ void vadd_hip(hipLaunchParm lp, const float* a, const float* b, float* c, int N) {
int idx = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
if (idx < N) {
@@ -32,16 +31,15 @@ __global__ void vadd_hip(hipLaunchParm lp, const float *a, const float *b, float
}
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
int sizeElements = 1000000;
size_t sizeBytes = sizeElements * sizeof(float);
bool pass = true;
// Allocate host memory
float *A_h = (float*)malloc(sizeBytes);
float *B_h = (float*)malloc(sizeBytes);
float *C_h = (float*)malloc(sizeBytes);
float* A_h = (float*)malloc(sizeBytes);
float* B_h = (float*)malloc(sizeBytes);
float* C_h = (float*)malloc(sizeBytes);
// Allocate device memory:
float *A_d, *B_d, *C_d;
@@ -50,8 +48,8 @@ int main(int argc, char *argv[])
hipMalloc(&C_d, sizeBytes);
// Initialize host memory
for (int i=0; i<sizeElements; i++) {
A_h[i] = 1.618f * i;
for (int i = 0; i < sizeElements; i++) {
A_h[i] = 1.618f * i;
B_h[i] = 3.142f * i;
}
@@ -60,20 +58,20 @@ int main(int argc, char *argv[])
hipMemcpy(B_d, B_h, sizeBytes, hipMemcpyHostToDevice);
// Launch kernel onto default accelerator
int blockSize = 256; // pick arbitrary block size
int blocks = (sizeElements+blockSize-1)/blockSize; // round up to launch enough blocks
int blockSize = 256; // pick arbitrary block size
int blocks = (sizeElements + blockSize - 1) / blockSize; // round up to launch enough blocks
hipLaunchKernel(vadd_hip, dim3(blocks), dim3(blockSize), 0, 0, A_d, B_d, C_d, sizeElements);
// D2H Copy
hipMemcpy(C_h, C_d, sizeBytes, hipMemcpyDeviceToHost);
// Verify
for (int i=0; i<sizeElements; i++) {
float ref= 1.618f * i + 3.142f * i;
for (int i = 0; i < sizeElements; i++) {
float ref = 1.618f * i + 3.142f * i;
if (C_h[i] != ref) {
printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
pass = false;
}
};
if (pass) printf ("PASSED!\n");
if (pass) printf("PASSED!\n");
}
projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp
+10 -10
Zobrazit soubor
@@ -22,25 +22,25 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include<iostream>
#include<fstream>
#include<vector>
#include <iostream>
#include <fstream>
#include <vector>
#define LEN 64
#define SIZE LEN<<2
#define SIZE LEN << 2
#define fileName "test.co"
#define kernel_name "vadd"
int main(){
int main() {
float *A, *B, *C;
hipDeviceptr_t Ad, Bd, Cd;
A = new float[LEN];
B = new float[LEN];
C = new float[LEN];
for(uint32_t i=0;i<LEN;i++){
A[i] = i*1.0f;
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 1.0f;
C[i] = 0.0f;
}
@@ -65,16 +65,16 @@ int main(){
hipModuleGetFunction(&Function, Module, kernel_name);
int n = LEN;
void * args[4] = {&Ad, &Bd, &Cd, &n};
void* args[4] = {&Ad, &Bd, &Cd, &n};
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
hipMemcpyDtoH(C, Cd, SIZE);
int mismatchCount = 0;
for(uint32_t i=0;i<LEN;i++){
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] + B[i] != C[i]) {
mismatchCount++;
std::cout<<"error: mismatch " << A[i]<<" + "<<B[i]<<" != "<<C[i]<<std::endl;
std::cout << "error: mismatch " << A[i] << " + " << B[i] << " != " << C[i] << std::endl;
}
}
projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp
+19 -20
Zobrazit soubor
@@ -21,7 +21,6 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include <iostream>
@@ -33,22 +32,25 @@ THE SOFTWARE.
#endif
#define LEN 64
#define SIZE LEN<<2
#define SIZE LEN << 2
#define fileName "vcpy_kernel.code.adipose"
#define kernel_name "hello_world"
#define HIP_CHECK(status) \
if(status != hipSuccess) {std::cout<<"Got Status: "<<status<<" at Line: "<<__LINE__<<std::endl;exit(0);}
#define HIP_CHECK(status) \
if (status != hipSuccess) { \
std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \
exit(0); \
}
int main(){
int main() {
float *A, *B;
hipDeviceptr_t Ad, Bd;
A = new float[LEN];
B = new float[LEN];
for(uint32_t i=0;i<LEN;i++){
A[i] = i*1.0f;
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
@@ -68,36 +70,33 @@ int main(){
HIP_CHECK(hipModuleLoad(&Module, fileName));
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
uint32_t len = LEN;
uint32_t one = 1;
uint32_t len = LEN;
uint32_t one = 1;
struct {
void * _Ad;
void * _Bd;
void* _Ad;
void* _Bd;
} args;
args._Ad = Ad;
args._Bd = Bd;
size_t size = sizeof(args);
void *config[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIP_CHECK(hipHccModuleLaunchKernel(Function, LEN, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
HIP_CHECK(
hipHccModuleLaunchKernel(Function, LEN, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
hipMemcpyDtoH(B, Bd, SIZE);
int mismatchCount = 0;
for(uint32_t i=0;i<LEN;i++){
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] != B[i]) {
mismatchCount++;
std::cout<<"error: mismatch " << A[i]<<" != "<<B[i]<<std::endl;
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
}
}
projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp
+19 -19
Zobrazit soubor
@@ -28,22 +28,25 @@ THE SOFTWARE.
#include <hip/hip_hcc.h>
#define LEN 64
#define SIZE LEN<<2
#define SIZE LEN << 2
#define fileName "vcpy_kernel.code.adipose"
#define kernel_name "hello_world"
#define HIP_CHECK(status) \
if(status != hipSuccess) {std::cout<<"Got Status: "<<status<<" at Line: "<<__LINE__<<std::endl;exit(0);}
#define HIP_CHECK(status) \
if (status != hipSuccess) { \
std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \
exit(0); \
}
int main(){
int main() {
float *A, *B;
hipDeviceptr_t Ad, Bd;
A = new float[LEN];
B = new float[LEN];
for(uint32_t i=0;i<LEN;i++){
A[i] = i*1.0f;
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
@@ -64,12 +67,12 @@ int main(){
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
#ifdef __HIP_PLATFORM_HCC__
uint32_t len = LEN;
uint32_t one = 1;
uint32_t len = LEN;
uint32_t one = 1;
struct {
void * _Ad;
void * _Bd;
void* _Ad;
void* _Bd;
} args;
args._Ad = Ad;
@@ -80,8 +83,8 @@ int main(){
#ifdef __HIP_PLATFORM_NVCC__
struct {
uint32_t _hidden[1];
void * _Ad;
void * _Bd;
void* _Ad;
void* _Bd;
} args;
args._hidden[0] = 0;
@@ -92,21 +95,18 @@ int main(){
size_t size = sizeof(args);
void *config[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
hipMemcpyDtoH(B, Bd, SIZE);
int mismatchCount = 0;
for(uint32_t i=0;i<LEN;i++){
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] != B[i]) {
mismatchCount++;
std::cout<<"error: mismatch " << A[i]<<" != "<<B[i]<<std::endl;
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
}
}
projects/hip-tests/samples/0_Intro/module_api/vcpy_kernel.cpp
+1 -2
Zobrazit soubor
@@ -22,8 +22,7 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
extern "C" __global__ void hello_world(float *a, float *b)
{
extern "C" __global__ void hello_world(float* a, float* b) {
int tx = hipThreadIdx_x;
b[tx] = a[tx];
}
projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp
+30 -29
Zobrazit soubor
@@ -28,25 +28,29 @@ THE SOFTWARE.
#include <hip/hip_hcc.h>
#define LEN 64
#define SIZE LEN*sizeof(float)
#define SIZE LEN * sizeof(float)
#define fileName "vcpy_kernel.code.adipose"
float myDeviceGlobal;
float myDeviceGlobalArray[16];
#define HIP_CHECK(cmd) \
{\
hipError_t status = cmd;\
if(status != hipSuccess) {std::cout<<"error: #"<<status<<" ("<< hipGetErrorString(status) << ") at line:"<<__LINE__<<": "<<#cmd<<std::endl;abort();}\
}
#define HIP_CHECK(cmd) \
{ \
hipError_t status = cmd; \
if (status != hipSuccess) { \
std::cout << "error: #" << status << " (" << hipGetErrorString(status) \
<< ") at line:" << __LINE__ << ": " << #cmd << std::endl; \
abort(); \
} \
}
int main(){
int main() {
float *A, *B;
float* Ad, *Bd;
float *Ad, *Bd;
A = new float[LEN];
B = new float[LEN];
for(uint32_t i=0;i<LEN;i++){
A[i] = i*1.0f;
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
@@ -70,18 +74,18 @@ int main(){
#define ARRAY_SIZE 16
float myDeviceGlobalArray_h[ARRAY_SIZE];
for (int i=0; i<ARRAY_SIZE; i++) {
myDeviceGlobalArray_h[i] = i*1000.0f;
myDeviceGlobalArray[i] = i*1000.0f;
for (int i = 0; i < ARRAY_SIZE; i++) {
myDeviceGlobalArray_h[i] = i * 1000.0f;
myDeviceGlobalArray[i] = i * 1000.0f;
}
#ifdef __HIP_PLATFORM_HCC__
uint32_t len = LEN;
uint32_t one = 1;
uint32_t len = LEN;
uint32_t one = 1;
struct {
void * _Ad;
void * _Bd;
void* _Ad;
void* _Bd;
} args;
args._Ad = Ad;
@@ -92,8 +96,8 @@ int main(){
#ifdef __HIP_PLATFORM_NVCC__
struct {
uint32_t _hidden[1];
void * _Ad;
void * _Bd;
void* _Ad;
void* _Bd;
} args;
args._hidden[0] = 0;
@@ -104,11 +108,8 @@ int main(){
size_t size = sizeof(args);
void *config[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
{
hipFunction_t Function;
@@ -118,10 +119,10 @@ int main(){
hipMemcpyDtoH(B, Bd, SIZE);
int mismatchCount = 0;
for(uint32_t i=0;i<LEN;i++){
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] != B[i]) {
mismatchCount++;
std::cout<<"error: mismatch " << A[i]<<" != "<<B[i]<<std::endl;
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
if (mismatchCount >= 10) {
break;
}
@@ -143,11 +144,11 @@ int main(){
hipMemcpyDtoH(B, Bd, SIZE);
int mismatchCount = 0;
for(uint32_t i=0;i<LEN;i++){
float expected = A[i] + myDeviceGlobal_h + myDeviceGlobalArray_h[i%16];
for (uint32_t i = 0; i < LEN; i++) {
float expected = A[i] + myDeviceGlobal_h + myDeviceGlobalArray_h[i % 16];
if (expected != B[i]) {
mismatchCount++;
std::cout<<"error: mismatch " << expected <<" != "<<B[i]<<std::endl;
std::cout << "error: mismatch " << expected << " != " << B[i] << std::endl;
if (mismatchCount >= 10) {
break;
}
projects/hip-tests/samples/0_Intro/module_api_global/vcpy_kernel.cpp
+5 -7
Zobrazit soubor
@@ -25,17 +25,15 @@ THE SOFTWARE.
#define ARRAY_SIZE (16)
extern float myDeviceGlobal;
extern float myDeviceGlobalArray[16];;
extern float myDeviceGlobalArray[16];
;
extern "C" __global__ void hello_world(const float *a, float *b)
{
extern "C" __global__ void hello_world(const float* a, float* b) {
int tx = hipThreadIdx_x;
b[tx] = a[tx];
}
extern "C" __global__ void test_globals(const float *a, float *b)
{
extern "C" __global__ void test_globals(const float* a, float* b) {
int tx = hipThreadIdx_x;
b[tx] = a[tx] + myDeviceGlobal+ myDeviceGlobalArray[tx%ARRAY_SIZE] ;
b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE];
}
projects/hip-tests/samples/0_Intro/square/square.hipref.cpp
+32 -35
Zobrazit soubor
@@ -23,33 +23,31 @@ THE SOFTWARE.
#include <stdio.h>
#include "hip/hip_runtime.h"
#define CHECK(cmd) \
{\
hipError_t error = cmd;\
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
exit(EXIT_FAILURE);\
}\
}
#define CHECK(cmd) \
{ \
hipError_t error = cmd; \
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \
__FILE__, __LINE__); \
exit(EXIT_FAILURE); \
} \
}
/*
* Square each element in the array A and write to array C.
*/
template <typename T>
__global__ void
vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N)
{
__global__ void vector_square(hipLaunchParm lp, T* C_d, const T* A_d, size_t N) {
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x ;
size_t stride = hipBlockDim_x * hipGridDim_x;
for (size_t i=offset; i<N; i+=stride) {
for (size_t i = offset; i < N; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
float *A_d, *C_d;
float *A_h, *C_h;
size_t N = 1000000;
@@ -57,43 +55,42 @@ int main(int argc, char *argv[])
static int device = 0;
CHECK(hipSetDevice(device));
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, device/*deviceID*/));
printf ("info: running on device %s\n", props.name);
#ifdef __HIP_PLATFORM_HCC__
printf ("info: architecture on AMD GPU device is: %d\n",props.gcnArch);
#endif
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
printf("info: running on device %s\n", props.name);
#ifdef __HIP_PLATFORM_HCC__
printf("info: architecture on AMD GPU device is: %d\n", props.gcnArch);
#endif
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
// Fill with Phi + i
for (size_t i=0; i<N; i++)
{
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
}
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
printf ("info: copy Host2Device\n");
CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
printf("info: copy Host2Device\n");
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
printf ("info: launch 'vector_square' kernel\n");
printf("info: launch 'vector_square' kernel\n");
hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: copy Device2Host\n");
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf("info: copy Device2Host\n");
CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf ("info: check result\n");
for (size_t i=0; i<N; i++) {
printf("info: check result\n");
for (size_t i = 0; i < N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
}
}
printf ("PASSED!\n");
printf("PASSED!\n");
}
projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp
+98 -171
Zobrazit soubor
@@ -10,96 +10,72 @@ using namespace std;
#define SORT_RETAIN_ATTS_ORDER 1
bool ResultDatabase::Result::operator<(const Result &rhs) const
{
if (test < rhs.test)
return true;
if (test > rhs.test)
return false;
#if (SORT_RETAIN_ATTS_ORDER == 0)
bool ResultDatabase::Result::operator<(const Result& rhs) const {
if (test < rhs.test) return true;
if (test > rhs.test) return false;
#if (SORT_RETAIN_ATTS_ORDER == 0)
// For ties, sort by the value of the attribute:
if (atts < rhs.atts)
return true;
if (atts > rhs.atts)
return false;
if (atts < rhs.atts) return true;
if (atts > rhs.atts) return false;
#endif
return false; // less-operator returns false on equal
return false; // less-operator returns false on equal
}
double ResultDatabase::Result::GetMin() const
{
double ResultDatabase::Result::GetMin() const {
double r = FLT_MAX;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r = min(r, value[i]);
}
return r;
}
double ResultDatabase::Result::GetMax() const
{
double ResultDatabase::Result::GetMax() const {
double r = -FLT_MAX;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r = max(r, value[i]);
}
return r;
}
double ResultDatabase::Result::GetMedian() const
{
return GetPercentile(50);
}
double ResultDatabase::Result::GetMedian() const { return GetPercentile(50); }
double ResultDatabase::Result::GetPercentile(double q) const
{
double ResultDatabase::Result::GetPercentile(double q) const {
int n = value.size();
if (n == 0)
return FLT_MAX;
if (n == 1)
return value[0];
if (n == 0) return FLT_MAX;
if (n == 1) return value[0];
if (q <= 0)
return value[0];
if (q >= 100)
return value[n-1];
if (q <= 0) return value[0];
if (q >= 100) return value[n - 1];
double index = ((n + 1.) * q / 100.) - 1;
vector<double> sorted = value;
sort(sorted.begin(), sorted.end());
if (n == 2)
return (sorted[0] * (1 - q/100.) + sorted[1] * (q/100.));
if (n == 2) return (sorted[0] * (1 - q / 100.) + sorted[1] * (q / 100.));
int index_lo = int(index);
double frac = index - index_lo;
if (frac == 0)
return sorted[index_lo];
if (frac == 0) return sorted[index_lo];
double lo = sorted[index_lo];
double hi = sorted[index_lo + 1];
return lo + (hi-lo)*frac;
return lo + (hi - lo) * frac;
}
double ResultDatabase::Result::GetMean() const
{
double ResultDatabase::Result::GetMean() const {
double r = 0;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r += value[i];
}
return r / double(value.size());
}
double ResultDatabase::Result::GetStdDev() const
{
double ResultDatabase::Result::GetStdDev() const {
double r = 0;
double u = GetMean();
if (u == FLT_MAX)
return FLT_MAX;
for (int i=0; i<value.size(); i++)
{
if (u == FLT_MAX) return FLT_MAX;
for (int i = 0; i < value.size(); i++) {
r += (value[i] - u) * (value[i] - u);
}
r = sqrt(r / value.size());
@@ -107,58 +83,42 @@ double ResultDatabase::Result::GetStdDev() const
}
void ResultDatabase::AddResults(const string &test,
const string &atts,
const string &unit,
const vector<double> &values)
{
for (int i=0; i<values.size(); i++)
{
void ResultDatabase::AddResults(const string& test, const string& atts, const string& unit,
const vector<double>& values) {
for (int i = 0; i < values.size(); i++) {
AddResult(test, atts, unit, values[i]);
}
}
static string RemoveAllButLeadingSpaces(const string &a)
{
static string RemoveAllButLeadingSpaces(const string& a) {
string b;
int n = a.length();
int i = 0;
while (i<n && a[i] == ' ')
{
while (i < n && a[i] == ' ') {
b += a[i];
++i;
}
for (; i<n; i++)
{
if (a[i] != ' ' && a[i] != '\t')
b += a[i];
for (; i < n; i++) {
if (a[i] != ' ' && a[i] != '\t') b += a[i];
}
return b;
}
void ResultDatabase::AddResult(const string &test_orig,
const string &atts_orig,
const string &unit_orig,
double value)
{
void ResultDatabase::AddResult(const string& test_orig, const string& atts_orig,
const string& unit_orig, double value) {
string test = RemoveAllButLeadingSpaces(test_orig);
string atts = RemoveAllButLeadingSpaces(atts_orig);
string unit = RemoveAllButLeadingSpaces(unit_orig);
int index;
for (index = 0; index < results.size(); index++)
{
if (results[index].test == test &&
results[index].atts == atts)
{
if (results[index].unit != unit)
throw "Internal error: mixed units";
for (index = 0; index < results.size(); index++) {
if (results[index].test == test && results[index].atts == atts) {
if (results[index].unit != unit) throw "Internal error: mixed units";
break;
}
}
if (index >= results.size())
{
if (index >= results.size()) {
Result r;
r.test = test;
r.atts = atts;
@@ -192,41 +152,33 @@ void ResultDatabase::AddResult(const string &test_orig,
// Changed note about missing values to be worded a little better.
//
// ****************************************************************************
void ResultDatabase::DumpDetailed(ostream &out)
{
void ResultDatabase::DumpDetailed(ostream& out) {
vector<Result> sorted(results);
stable_sort(sorted.begin(), sorted.end());
const int testNameW = 24 ;
const int testNameW = 24;
const int attW = 12;
const int fieldW = 11;
out << std::fixed << right << std::setprecision(4);
int maxtrials = 1;
for (int i=0; i<sorted.size(); i++)
{
if (sorted[i].value.size() > maxtrials)
maxtrials = sorted[i].value.size();
for (int i = 0; i < sorted.size(); i++) {
if (sorted[i].value.size() > maxtrials) maxtrials = sorted[i].value.size();
}
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << setw(testNameW) << "test\t"
<< setw(attW) << "atts\t"
<< setw(fieldW)
<< "median\t"
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t"
<< "mean\t"
<< "stddev\t"
<< "min\t"
<< "max\t";
for (int i=0; i<maxtrials; i++)
out << "trial"<<i<<"\t";
for (int i = 0; i < maxtrials; i++) out << "trial" << i << "\t";
out << endl;
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << setw(testNameW) << r.test + "\t";
out << setw(attW) << r.atts + "\t";
out << setw(fieldW) << r.unit + "\t";
@@ -237,7 +189,7 @@ void ResultDatabase::DumpDetailed(ostream &out)
if (r.GetMean() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMean() << "\t";
out << r.GetMean() << "\t";
if (r.GetStdDev() == FLT_MAX)
out << "N/A\t";
else
@@ -245,13 +197,12 @@ void ResultDatabase::DumpDetailed(ostream &out)
if (r.GetMin() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMin() << "\t";
out << r.GetMin() << "\t";
if (r.GetMax() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMax() << "\t";
for (int j=0; j<r.value.size(); j++)
{
out << r.GetMax() << "\t";
for (int j = 0; j < r.value.size(); j++) {
if (r.value[j] == FLT_MAX)
out << "N/A\t";
else
@@ -285,23 +236,19 @@ void ResultDatabase::DumpDetailed(ostream &out)
// Added note about (*) missing value tag.
//
// ****************************************************************************
void ResultDatabase::DumpSummary(ostream &out)
{
void ResultDatabase::DumpSummary(ostream& out) {
vector<Result> sorted(results);
stable_sort(sorted.begin(), sorted.end());
const int testNameW = 24 ;
const int testNameW = 24;
const int attW = 12;
const int fieldW = 9;
out << std::fixed << right << std::setprecision(4);
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << setw(testNameW) << "test\t"
<< setw(attW) << "atts\t"
<< setw(fieldW)
<< "units\t"
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t"
<< "median\t"
<< "mean\t"
<< "stddev\t"
@@ -309,9 +256,8 @@ void ResultDatabase::DumpSummary(ostream &out)
<< "max\t";
out << endl;
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << setw(testNameW) << r.test + "\t";
out << setw(attW) << r.atts + "\t";
out << setw(fieldW) << r.unit + "\t";
@@ -322,7 +268,7 @@ void ResultDatabase::DumpSummary(ostream &out)
if (r.GetMean() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMean() << "\t";
out << r.GetMean() << "\t";
if (r.GetStdDev() == FLT_MAX)
out << "N/A\t";
else
@@ -330,11 +276,11 @@ void ResultDatabase::DumpSummary(ostream &out)
if (r.GetMin() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMin() << "\t";
out << r.GetMin() << "\t";
if (r.GetMax() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMax() << "\t";
out << r.GetMax() << "\t";
out << endl;
}
@@ -359,10 +305,7 @@ void ResultDatabase::DumpSummary(ostream &out)
//
//
// ****************************************************************************
void ResultDatabase::ClearAllResults()
{
results.clear();
}
void ResultDatabase::ClearAllResults() { results.clear(); }
// ****************************************************************************
// Method: ResultDatabase::DumpCsv
@@ -380,39 +323,36 @@ void ResultDatabase::ClearAllResults()
// Modifications:
//
// ****************************************************************************
void ResultDatabase::DumpCsv(string fileName)
{
void ResultDatabase::DumpCsv(string fileName) {
bool emptyFile;
vector<Result> sorted(results);
stable_sort(sorted.begin(), sorted.end());
//Check to see if the file is empty - if so, add the headers
// Check to see if the file is empty - if so, add the headers
emptyFile = this->IsFileEmpty(fileName);
//Open file and append by default
// Open file and append by default
ofstream out;
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
//Add headers only for empty files
if(emptyFile)
{
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << "test, "
<< "atts, "
<< "units, "
<< "median, "
<< "mean, "
<< "stddev, "
<< "min, "
<< "max, ";
out << endl;
// Add headers only for empty files
if (emptyFile) {
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << "test, "
<< "atts, "
<< "units, "
<< "median, "
<< "mean, "
<< "stddev, "
<< "min, "
<< "max, ";
out << endl;
}
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << r.test << ", ";
out << r.atts << ", ";
out << r.unit << ", ";
@@ -423,7 +363,7 @@ void ResultDatabase::DumpCsv(string fileName)
if (r.GetMean() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMean() << ", ";
out << r.GetMean() << ", ";
if (r.GetStdDev() == FLT_MAX)
out << "N/A, ";
else
@@ -431,11 +371,11 @@ void ResultDatabase::DumpCsv(string fileName)
if (r.GetMin() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMin() << ", ";
out << r.GetMin() << ", ";
if (r.GetMax() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMax() << ", ";
out << r.GetMax() << ", ";
out << endl;
}
@@ -460,29 +400,24 @@ void ResultDatabase::DumpCsv(string fileName)
//
// ****************************************************************************
bool ResultDatabase::IsFileEmpty(string fileName)
{
bool fileEmpty;
bool ResultDatabase::IsFileEmpty(string fileName) {
bool fileEmpty;
ifstream file(fileName.c_str());
ifstream file(fileName.c_str());
//If the file doesn't exist it is by definition empty
if(!file.good())
{
// If the file doesn't exist it is by definition empty
if (!file.good()) {
return true;
}
else
{
} else {
fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
file.close();
return fileEmpty;
}
//Otherwise, return false
return false;
}
return fileEmpty;
}
// Otherwise, return false
return false;
}
// ****************************************************************************
@@ -500,16 +435,12 @@ bool ResultDatabase::IsFileEmpty(string fileName)
// Modifications:
//
// ****************************************************************************
vector<ResultDatabase::Result>
ResultDatabase::GetResultsForTest(const string &test)
{
vector<ResultDatabase::Result> ResultDatabase::GetResultsForTest(const string& test) {
// get only the given test results
vector<Result> retval;
for (int i=0; i<results.size(); i++)
{
Result &r = results[i];
if (r.test == test)
retval.push_back(r);
for (int i = 0; i < results.size(); i++) {
Result& r = results[i];
if (r.test == test) retval.push_back(r);
}
return retval;
}
@@ -528,8 +459,4 @@ ResultDatabase::GetResultsForTest(const string &test)
// Modifications:
//
// ****************************************************************************
const vector<ResultDatabase::Result> &
ResultDatabase::GetResults() const
{
return results;
}
const vector<ResultDatabase::Result>& ResultDatabase::GetResults() const { return results; }
projects/hip-tests/samples/1_Utils/hipBusBandwidth/ResultDatabase.h
+22 -33
Zobrazit soubor
@@ -6,11 +6,11 @@
#include <iostream>
#include <fstream>
#include <cfloat>
using std::ifstream;
using std::ofstream;
using std::ostream;
using std::string;
using std::vector;
using std::ostream;
using std::ofstream;
using std::ifstream;
// ****************************************************************************
@@ -40,18 +40,16 @@ using std::ifstream;
// Added a GetResults method as well, and made several functions const.
//
// ****************************************************************************
class ResultDatabase
{
public:
class ResultDatabase {
public:
//
// A performance result for a single SHOC benchmark run.
//
struct Result
{
string test; // e.g. "readback"
string atts; // e.g. "pagelocked 4k^2"
string unit; // e.g. "MB/sec"
vector<double> value; // e.g. "837.14"
struct Result {
string test; // e.g. "readback"
string atts; // e.g. "pagelocked 4k^2"
string unit; // e.g. "MB/sec"
vector<double> value; // e.g. "837.14"
double GetMin() const;
double GetMax() const;
double GetMedian() const;
@@ -59,41 +57,32 @@ class ResultDatabase
double GetMean() const;
double GetStdDev() const;
bool operator<(const Result &rhs) const;
bool operator<(const Result& rhs) const;
bool HadAnyFLTMAXValues() const
{
for (int i=0; i<value.size(); ++i)
{
if (value[i] >= FLT_MAX)
return true;
bool HadAnyFLTMAXValues() const {
for (int i = 0; i < value.size(); ++i) {
if (value[i] >= FLT_MAX) return true;
}
return false;
}
};
protected:
protected:
vector<Result> results;
public:
void AddResult(const string &test,
const string &atts,
const string &unit,
double value);
void AddResults(const string &test,
const string &atts,
const string &unit,
const vector<double> &values);
vector<Result> GetResultsForTest(const string &test);
const vector<Result> &GetResults() const;
public:
void AddResult(const string& test, const string& atts, const string& unit, double value);
void AddResults(const string& test, const string& atts, const string& unit,
const vector<double>& values);
vector<Result> GetResultsForTest(const string& test);
const vector<Result>& GetResults() const;
void ClearAllResults();
void DumpDetailed(ostream&);
void DumpSummary(ostream&);
void DumpCsv(string fileName);
private:
private:
bool IsFileEmpty(string fileName);
};
projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp
+329 -376
Zobrazit soubor
Rozdílový obsah nebyl zobrazen, protože je příliš veliký Načíst rozdílové porovnání
projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.cpp
+97 -170
Zobrazit soubor
@@ -7,93 +7,69 @@
using namespace std;
bool ResultDatabase::Result::operator<(const Result &rhs) const
{
if (test < rhs.test)
return true;
if (test > rhs.test)
return false;
if (atts < rhs.atts)
return true;
if (atts > rhs.atts)
return false;
return false; // less-operator returns false on equal
bool ResultDatabase::Result::operator<(const Result& rhs) const {
if (test < rhs.test) return true;
if (test > rhs.test) return false;
if (atts < rhs.atts) return true;
if (atts > rhs.atts) return false;
return false; // less-operator returns false on equal
}
double ResultDatabase::Result::GetMin() const
{
double ResultDatabase::Result::GetMin() const {
double r = FLT_MAX;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r = min(r, value[i]);
}
return r;
}
double ResultDatabase::Result::GetMax() const
{
double ResultDatabase::Result::GetMax() const {
double r = -FLT_MAX;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r = max(r, value[i]);
}
return r;
}
double ResultDatabase::Result::GetMedian() const
{
return GetPercentile(50);
}
double ResultDatabase::Result::GetMedian() const { return GetPercentile(50); }
double ResultDatabase::Result::GetPercentile(double q) const
{
double ResultDatabase::Result::GetPercentile(double q) const {
int n = value.size();
if (n == 0)
return FLT_MAX;
if (n == 1)
return value[0];
if (n == 0) return FLT_MAX;
if (n == 1) return value[0];
if (q <= 0)
return value[0];
if (q >= 100)
return value[n-1];
if (q <= 0) return value[0];
if (q >= 100) return value[n - 1];
double index = ((n + 1.) * q / 100.) - 1;
vector<double> sorted = value;
sort(sorted.begin(), sorted.end());
if (n == 2)
return (sorted[0] * (1 - q/100.) + sorted[1] * (q/100.));
if (n == 2) return (sorted[0] * (1 - q / 100.) + sorted[1] * (q / 100.));
int index_lo = int(index);
double frac = index - index_lo;
if (frac == 0)
return sorted[index_lo];
if (frac == 0) return sorted[index_lo];
double lo = sorted[index_lo];
double hi = sorted[index_lo + 1];
return lo + (hi-lo)*frac;
return lo + (hi - lo) * frac;
}
double ResultDatabase::Result::GetMean() const
{
double ResultDatabase::Result::GetMean() const {
double r = 0;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r += value[i];
}
return r / double(value.size());
}
double ResultDatabase::Result::GetStdDev() const
{
double ResultDatabase::Result::GetStdDev() const {
double r = 0;
double u = GetMean();
if (u == FLT_MAX)
return FLT_MAX;
for (int i=0; i<value.size(); i++)
{
if (u == FLT_MAX) return FLT_MAX;
for (int i = 0; i < value.size(); i++) {
r += (value[i] - u) * (value[i] - u);
}
r = sqrt(r / value.size());
@@ -101,58 +77,42 @@ double ResultDatabase::Result::GetStdDev() const
}
void ResultDatabase::AddResults(const string &test,
const string &atts,
const string &unit,
const vector<double> &values)
{
for (int i=0; i<values.size(); i++)
{
void ResultDatabase::AddResults(const string& test, const string& atts, const string& unit,
const vector<double>& values) {
for (int i = 0; i < values.size(); i++) {
AddResult(test, atts, unit, values[i]);
}
}
static string RemoveAllButLeadingSpaces(const string &a)
{
static string RemoveAllButLeadingSpaces(const string& a) {
string b;
int n = a.length();
int i = 0;
while (i<n && a[i] == ' ')
{
while (i < n && a[i] == ' ') {
b += a[i];
++i;
}
for (; i<n; i++)
{
if (a[i] != ' ' && a[i] != '\t')
b += a[i];
for (; i < n; i++) {
if (a[i] != ' ' && a[i] != '\t') b += a[i];
}
return b;
}
void ResultDatabase::AddResult(const string &test_orig,
const string &atts_orig,
const string &unit_orig,
double value)
{
void ResultDatabase::AddResult(const string& test_orig, const string& atts_orig,
const string& unit_orig, double value) {
string test = RemoveAllButLeadingSpaces(test_orig);
string atts = RemoveAllButLeadingSpaces(atts_orig);
string unit = RemoveAllButLeadingSpaces(unit_orig);
int index;
for (index = 0; index < results.size(); index++)
{
if (results[index].test == test &&
results[index].atts == atts)
{
if (results[index].unit != unit)
throw "Internal error: mixed units";
for (index = 0; index < results.size(); index++) {
if (results[index].test == test && results[index].atts == atts) {
if (results[index].unit != unit) throw "Internal error: mixed units";
break;
}
}
if (index >= results.size())
{
if (index >= results.size()) {
Result r;
r.test = test;
r.atts = atts;
@@ -186,40 +146,32 @@ void ResultDatabase::AddResult(const string &test_orig,
// Changed note about missing values to be worded a little better.
//
// ****************************************************************************
void ResultDatabase::DumpDetailed(ostream &out)
{
void ResultDatabase::DumpDetailed(ostream& out) {
vector<Result> sorted(results);
sort(sorted.begin(), sorted.end());
const int testNameW = 24 ;
const int testNameW = 24;
const int attW = 12;
const int fieldW = 11;
out << std::fixed << right << std::setprecision(4);
int maxtrials = 1;
for (int i=0; i<sorted.size(); i++)
{
if (sorted[i].value.size() > maxtrials)
maxtrials = sorted[i].value.size();
for (int i = 0; i < sorted.size(); i++) {
if (sorted[i].value.size() > maxtrials) maxtrials = sorted[i].value.size();
}
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << setw(testNameW) << "test\t"
<< setw(attW) << "atts\t"
<< setw(fieldW)
<< "median\t"
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t"
<< "mean\t"
<< "stddev\t"
<< "min\t"
<< "max\t";
for (int i=0; i<maxtrials; i++)
out << "trial"<<i<<"\t";
for (int i = 0; i < maxtrials; i++) out << "trial" << i << "\t";
out << endl;
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << setw(testNameW) << r.test + "\t";
out << setw(attW) << r.atts + "\t";
out << setw(fieldW) << r.unit + "\t";
@@ -230,7 +182,7 @@ void ResultDatabase::DumpDetailed(ostream &out)
if (r.GetMean() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMean() << "\t";
out << r.GetMean() << "\t";
if (r.GetStdDev() == FLT_MAX)
out << "N/A\t";
else
@@ -238,13 +190,12 @@ void ResultDatabase::DumpDetailed(ostream &out)
if (r.GetMin() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMin() << "\t";
out << r.GetMin() << "\t";
if (r.GetMax() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMax() << "\t";
for (int j=0; j<r.value.size(); j++)
{
out << r.GetMax() << "\t";
for (int j = 0; j < r.value.size(); j++) {
if (r.value[j] == FLT_MAX)
out << "N/A\t";
else
@@ -278,22 +229,18 @@ void ResultDatabase::DumpDetailed(ostream &out)
// Added note about (*) missing value tag.
//
// ****************************************************************************
void ResultDatabase::DumpSummary(ostream &out)
{
void ResultDatabase::DumpSummary(ostream& out) {
vector<Result> sorted(results);
sort(sorted.begin(), sorted.end());
const int testNameW = 24 ;
const int testNameW = 24;
const int attW = 12;
const int fieldW = 9;
out << std::fixed << right << std::setprecision(4);
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << setw(testNameW) << "test\t"
<< setw(attW) << "atts\t"
<< setw(fieldW)
<< "units\t"
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t"
<< "median\t"
<< "mean\t"
<< "stddev\t"
@@ -301,9 +248,8 @@ void ResultDatabase::DumpSummary(ostream &out)
<< "max\t";
out << endl;
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << setw(testNameW) << r.test + "\t";
out << setw(attW) << r.atts + "\t";
out << setw(fieldW) << r.unit + "\t";
@@ -314,7 +260,7 @@ void ResultDatabase::DumpSummary(ostream &out)
if (r.GetMean() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMean() << "\t";
out << r.GetMean() << "\t";
if (r.GetStdDev() == FLT_MAX)
out << "N/A\t";
else
@@ -322,11 +268,11 @@ void ResultDatabase::DumpSummary(ostream &out)
if (r.GetMin() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMin() << "\t";
out << r.GetMin() << "\t";
if (r.GetMax() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMax() << "\t";
out << r.GetMax() << "\t";
out << endl;
}
@@ -351,10 +297,7 @@ void ResultDatabase::DumpSummary(ostream &out)
//
//
// ****************************************************************************
void ResultDatabase::ClearAllResults()
{
results.clear();
}
void ResultDatabase::ClearAllResults() { results.clear(); }
// ****************************************************************************
// Method: ResultDatabase::DumpCsv
@@ -372,39 +315,36 @@ void ResultDatabase::ClearAllResults()
// Modifications:
//
// ****************************************************************************
void ResultDatabase::DumpCsv(string fileName)
{
void ResultDatabase::DumpCsv(string fileName) {
bool emptyFile;
vector<Result> sorted(results);
sort(sorted.begin(), sorted.end());
//Check to see if the file is empty - if so, add the headers
// Check to see if the file is empty - if so, add the headers
emptyFile = this->IsFileEmpty(fileName);
//Open file and append by default
// Open file and append by default
ofstream out;
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
//Add headers only for empty files
if(emptyFile)
{
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << "test, "
<< "atts, "
<< "units, "
<< "median, "
<< "mean, "
<< "stddev, "
<< "min, "
<< "max, ";
out << endl;
// Add headers only for empty files
if (emptyFile) {
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << "test, "
<< "atts, "
<< "units, "
<< "median, "
<< "mean, "
<< "stddev, "
<< "min, "
<< "max, ";
out << endl;
}
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << r.test << ", ";
out << r.atts << ", ";
out << r.unit << ", ";
@@ -415,7 +355,7 @@ void ResultDatabase::DumpCsv(string fileName)
if (r.GetMean() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMean() << ", ";
out << r.GetMean() << ", ";
if (r.GetStdDev() == FLT_MAX)
out << "N/A, ";
else
@@ -423,11 +363,11 @@ void ResultDatabase::DumpCsv(string fileName)
if (r.GetMin() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMin() << ", ";
out << r.GetMin() << ", ";
if (r.GetMax() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMax() << ", ";
out << r.GetMax() << ", ";
out << endl;
}
@@ -452,29 +392,24 @@ void ResultDatabase::DumpCsv(string fileName)
//
// ****************************************************************************
bool ResultDatabase::IsFileEmpty(string fileName)
{
bool fileEmpty;
bool ResultDatabase::IsFileEmpty(string fileName) {
bool fileEmpty;
ifstream file(fileName.c_str());
ifstream file(fileName.c_str());
//If the file doesn't exist it is by definition empty
if(!file.good())
{
// If the file doesn't exist it is by definition empty
if (!file.good()) {
return true;
}
else
{
} else {
fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
file.close();
return fileEmpty;
}
//Otherwise, return false
return false;
}
return fileEmpty;
}
// Otherwise, return false
return false;
}
// ****************************************************************************
@@ -492,16 +427,12 @@ bool ResultDatabase::IsFileEmpty(string fileName)
// Modifications:
//
// ****************************************************************************
vector<ResultDatabase::Result>
ResultDatabase::GetResultsForTest(const string &test)
{
vector<ResultDatabase::Result> ResultDatabase::GetResultsForTest(const string& test) {
// get only the given test results
vector<Result> retval;
for (int i=0; i<results.size(); i++)
{
Result &r = results[i];
if (r.test == test)
retval.push_back(r);
for (int i = 0; i < results.size(); i++) {
Result& r = results[i];
if (r.test == test) retval.push_back(r);
}
return retval;
}
@@ -520,8 +451,4 @@ ResultDatabase::GetResultsForTest(const string &test)
// Modifications:
//
// ****************************************************************************
const vector<ResultDatabase::Result> &
ResultDatabase::GetResults() const
{
return results;
}
const vector<ResultDatabase::Result>& ResultDatabase::GetResults() const { return results; }
projects/hip-tests/samples/1_Utils/hipCommander/ResultDatabase.h
+22 -33
Zobrazit soubor
@@ -6,11 +6,11 @@
#include <iostream>
#include <fstream>
#include <cfloat>
using std::ifstream;
using std::ofstream;
using std::ostream;
using std::string;
using std::vector;
using std::ostream;
using std::ofstream;
using std::ifstream;
// ****************************************************************************
@@ -40,18 +40,16 @@ using std::ifstream;
// Added a GetResults method as well, and made several functions const.
//
// ****************************************************************************
class ResultDatabase
{
public:
class ResultDatabase {
public:
//
// A performance result for a single SHOC benchmark run.
//
struct Result
{
string test; // e.g. "readback"
string atts; // e.g. "pagelocked 4k^2"
string unit; // e.g. "MB/sec"
vector<double> value; // e.g. "837.14"
struct Result {
string test; // e.g. "readback"
string atts; // e.g. "pagelocked 4k^2"
string unit; // e.g. "MB/sec"
vector<double> value; // e.g. "837.14"
double GetMin() const;
double GetMax() const;
double GetMedian() const;
@@ -59,41 +57,32 @@ class ResultDatabase
double GetMean() const;
double GetStdDev() const;
bool operator<(const Result &rhs) const;
bool operator<(const Result& rhs) const;
bool HadAnyFLTMAXValues() const
{
for (int i=0; i<value.size(); ++i)
{
if (value[i] >= FLT_MAX)
return true;
bool HadAnyFLTMAXValues() const {
for (int i = 0; i < value.size(); ++i) {
if (value[i] >= FLT_MAX) return true;
}
return false;
}
};
protected:
protected:
vector<Result> results;
public:
void AddResult(const string &test,
const string &atts,
const string &unit,
double value);
void AddResults(const string &test,
const string &atts,
const string &unit,
const vector<double> &values);
vector<Result> GetResultsForTest(const string &test);
const vector<Result> &GetResults() const;
public:
void AddResult(const string& test, const string& atts, const string& unit, double value);
void AddResults(const string& test, const string& atts, const string& unit,
const vector<double>& values);
vector<Result> GetResultsForTest(const string& test);
const vector<Result>& GetResults() const;
void ClearAllResults();
void DumpDetailed(ostream&);
void DumpSummary(ostream&);
void DumpCsv(string fileName);
private:
private:
bool IsFileEmpty(string fileName);
};
projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp
+324 -420
Zobrazit soubor
Rozdílový obsah nebyl zobrazen, protože je příliš veliký Načíst rozdílové porovnání
projects/hip-tests/samples/1_Utils/hipCommander/nullkernel.hip.cpp
+1 -1
Zobrazit soubor
@@ -1,6 +1,6 @@
#include "hip/hip_runtime.h"
extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad){
extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad) {
if (Ad) {
Ad[0] = 42;
}
projects/hip-tests/samples/1_Utils/hipCommander/testcase.cpp
+8 -11
Zobrazit soubor
@@ -1,20 +1,17 @@
#include <hip/hip_runtime.h>
static const int BLOCKSIZEX=32;
static const int BLOCKSIZEY=16;
static const int BLOCKSIZEX = 32;
static const int BLOCKSIZEY = 16;
__global__ void fails(hipLaunchParm lp, float* pErrorI)
{
if(pErrorI!=0)
{
pErrorI[0]=1;
__global__ void fails(hipLaunchParm lp, float* pErrorI) {
if (pErrorI != 0) {
pErrorI[0] = 1;
}
}
int main()
{
dim3 blocks(1,1);
dim3 threads(BLOCKSIZEX,BLOCKSIZEY);
int main() {
dim3 blocks(1, 1);
dim3 threads(BLOCKSIZEX, BLOCKSIZEY);
float error;
hipLaunchKernel(HIP_KERNEL_NAME(fails), blocks, threads, 0, 0, &error);
projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp
+98 -171
Zobrazit soubor
@@ -11,96 +11,72 @@ using namespace std;
#define SORT_RETAIN_ATTS_ORDER 1
bool ResultDatabase::Result::operator<(const Result &rhs) const
{
if (test < rhs.test)
return true;
if (test > rhs.test)
return false;
#if (SORT_RETAIN_ATTS_ORDER == 0)
bool ResultDatabase::Result::operator<(const Result& rhs) const {
if (test < rhs.test) return true;
if (test > rhs.test) return false;
#if (SORT_RETAIN_ATTS_ORDER == 0)
// For ties, sort by the value of the attribute:
if (atts < rhs.atts)
return true;
if (atts > rhs.atts)
return false;
if (atts < rhs.atts) return true;
if (atts > rhs.atts) return false;
#endif
return false; // less-operator returns false on equal
return false; // less-operator returns false on equal
}
double ResultDatabase::Result::GetMin() const
{
double ResultDatabase::Result::GetMin() const {
double r = FLT_MAX;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r = min(r, value[i]);
}
return r;
}
double ResultDatabase::Result::GetMax() const
{
double ResultDatabase::Result::GetMax() const {
double r = -FLT_MAX;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r = max(r, value[i]);
}
return r;
}
double ResultDatabase::Result::GetMedian() const
{
return GetPercentile(50);
}
double ResultDatabase::Result::GetMedian() const { return GetPercentile(50); }
double ResultDatabase::Result::GetPercentile(double q) const
{
double ResultDatabase::Result::GetPercentile(double q) const {
int n = value.size();
if (n == 0)
return FLT_MAX;
if (n == 1)
return value[0];
if (n == 0) return FLT_MAX;
if (n == 1) return value[0];
if (q <= 0)
return value[0];
if (q >= 100)
return value[n-1];
if (q <= 0) return value[0];
if (q >= 100) return value[n - 1];
double index = ((n + 1.) * q / 100.) - 1;
vector<double> sorted = value;
sort(sorted.begin(), sorted.end());
if (n == 2)
return (sorted[0] * (1 - q/100.) + sorted[1] * (q/100.));
if (n == 2) return (sorted[0] * (1 - q / 100.) + sorted[1] * (q / 100.));
int index_lo = int(index);
double frac = index - index_lo;
if (frac == 0)
return sorted[index_lo];
if (frac == 0) return sorted[index_lo];
double lo = sorted[index_lo];
double hi = sorted[index_lo + 1];
return lo + (hi-lo)*frac;
return lo + (hi - lo) * frac;
}
double ResultDatabase::Result::GetMean() const
{
double ResultDatabase::Result::GetMean() const {
double r = 0;
for (int i=0; i<value.size(); i++)
{
for (int i = 0; i < value.size(); i++) {
r += value[i];
}
return r / double(value.size());
}
double ResultDatabase::Result::GetStdDev() const
{
double ResultDatabase::Result::GetStdDev() const {
double r = 0;
double u = GetMean();
if (u == FLT_MAX)
return FLT_MAX;
for (int i=0; i<value.size(); i++)
{
if (u == FLT_MAX) return FLT_MAX;
for (int i = 0; i < value.size(); i++) {
r += (value[i] - u) * (value[i] - u);
}
r = sqrt(r / value.size());
@@ -108,58 +84,42 @@ double ResultDatabase::Result::GetStdDev() const
}
void ResultDatabase::AddResults(const string &test,
const string &atts,
const string &unit,
const vector<double> &values)
{
for (int i=0; i<values.size(); i++)
{
void ResultDatabase::AddResults(const string& test, const string& atts, const string& unit,
const vector<double>& values) {
for (int i = 0; i < values.size(); i++) {
AddResult(test, atts, unit, values[i]);
}
}
static string RemoveAllButLeadingSpaces(const string &a)
{
static string RemoveAllButLeadingSpaces(const string& a) {
string b;
int n = a.length();
int i = 0;
while (i<n && a[i] == ' ')
{
while (i < n && a[i] == ' ') {
b += a[i];
++i;
}
for (; i<n; i++)
{
if (a[i] != ' ' && a[i] != '\t')
b += a[i];
for (; i < n; i++) {
if (a[i] != ' ' && a[i] != '\t') b += a[i];
}
return b;
}
void ResultDatabase::AddResult(const string &test_orig,
const string &atts_orig,
const string &unit_orig,
double value)
{
void ResultDatabase::AddResult(const string& test_orig, const string& atts_orig,
const string& unit_orig, double value) {
string test = RemoveAllButLeadingSpaces(test_orig);
string atts = RemoveAllButLeadingSpaces(atts_orig);
string unit = RemoveAllButLeadingSpaces(unit_orig);
int index;
for (index = 0; index < results.size(); index++)
{
if (results[index].test == test &&
results[index].atts == atts)
{
if (results[index].unit != unit)
throw "Internal error: mixed units";
for (index = 0; index < results.size(); index++) {
if (results[index].test == test && results[index].atts == atts) {
if (results[index].unit != unit) throw "Internal error: mixed units";
break;
}
}
if (index >= results.size())
{
if (index >= results.size()) {
Result r;
r.test = test;
r.atts = atts;
@@ -193,43 +153,35 @@ void ResultDatabase::AddResult(const string &test_orig,
// Changed note about missing values to be worded a little better.
//
// ****************************************************************************
void ResultDatabase::DumpDetailed(ostream &out)
{
void ResultDatabase::DumpDetailed(ostream& out) {
vector<Result> sorted(results);
#if SORT_BY_NAME
stable_sort(sorted.begin(), sorted.end());
#endif
const int testNameW = 24 ;
const int testNameW = 24;
const int attW = 12;
const int fieldW = 11;
out << std::fixed << right << std::setprecision(4);
int maxtrials = 1;
for (int i=0; i<sorted.size(); i++)
{
if (sorted[i].value.size() > maxtrials)
maxtrials = sorted[i].value.size();
for (int i = 0; i < sorted.size(); i++) {
if (sorted[i].value.size() > maxtrials) maxtrials = sorted[i].value.size();
}
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << setw(testNameW) << "test\t"
<< setw(attW) << "atts\t"
<< setw(fieldW)
<< "median\t"
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t"
<< "mean\t"
<< "stddev\t"
<< "min\t"
<< "max\t";
for (int i=0; i<maxtrials; i++)
out << "trial"<<i<<"\t";
for (int i = 0; i < maxtrials; i++) out << "trial" << i << "\t";
out << endl;
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << setw(testNameW) << r.test + "\t";
out << setw(attW) << r.atts + "\t";
out << setw(fieldW) << r.unit + "\t";
@@ -240,7 +192,7 @@ void ResultDatabase::DumpDetailed(ostream &out)
if (r.GetMean() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMean() << "\t";
out << r.GetMean() << "\t";
if (r.GetStdDev() == FLT_MAX)
out << "N/A\t";
else
@@ -248,13 +200,12 @@ void ResultDatabase::DumpDetailed(ostream &out)
if (r.GetMin() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMin() << "\t";
out << r.GetMin() << "\t";
if (r.GetMax() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMax() << "\t";
for (int j=0; j<r.value.size(); j++)
{
out << r.GetMax() << "\t";
for (int j = 0; j < r.value.size(); j++) {
if (r.value[j] == FLT_MAX)
out << "N/A\t";
else
@@ -290,25 +241,21 @@ void ResultDatabase::DumpDetailed(ostream &out)
// Added note about (*) missing value tag.
//
// ****************************************************************************
void ResultDatabase::DumpSummary(ostream &out)
{
void ResultDatabase::DumpSummary(ostream& out) {
vector<Result> sorted(results);
#if SORT_BY_NAME
stable_sort(sorted.begin(), sorted.end());
#endif
const int testNameW = 32 ;
const int testNameW = 32;
const int attW = 12;
const int fieldW = 9;
out << std::fixed << right << std::setprecision(2);
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << setw(testNameW) << "test\t"
<< setw(attW) << "atts\t"
<< setw(fieldW)
<< "units\t"
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t"
<< "median\t"
<< "mean\t"
<< "stddev\t"
@@ -316,9 +263,8 @@ void ResultDatabase::DumpSummary(ostream &out)
<< "max\t";
out << endl;
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << setw(testNameW) << r.test + "\t";
out << setw(attW) << r.atts + "\t";
out << setw(fieldW) << r.unit + "\t";
@@ -329,7 +275,7 @@ void ResultDatabase::DumpSummary(ostream &out)
if (r.GetMean() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMean() << "\t";
out << r.GetMean() << "\t";
if (r.GetStdDev() == FLT_MAX)
out << "N/A\t";
else
@@ -337,11 +283,11 @@ void ResultDatabase::DumpSummary(ostream &out)
if (r.GetMin() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMin() << "\t";
out << r.GetMin() << "\t";
if (r.GetMax() == FLT_MAX)
out << "N/A\t";
else
out << r.GetMax() << "\t";
out << r.GetMax() << "\t";
out << endl;
}
@@ -368,10 +314,7 @@ void ResultDatabase::DumpSummary(ostream &out)
//
//
// ****************************************************************************
void ResultDatabase::ClearAllResults()
{
results.clear();
}
void ResultDatabase::ClearAllResults() { results.clear(); }
// ****************************************************************************
// Method: ResultDatabase::DumpCsv
@@ -389,8 +332,7 @@ void ResultDatabase::ClearAllResults()
// Modifications:
//
// ****************************************************************************
void ResultDatabase::DumpCsv(string fileName)
{
void ResultDatabase::DumpCsv(string fileName) {
bool emptyFile;
vector<Result> sorted(results);
@@ -398,32 +340,30 @@ void ResultDatabase::DumpCsv(string fileName)
stable_sort(sorted.begin(), sorted.end());
#endif
//Check to see if the file is empty - if so, add the headers
// Check to see if the file is empty - if so, add the headers
emptyFile = this->IsFileEmpty(fileName);
//Open file and append by default
// Open file and append by default
ofstream out;
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
//Add headers only for empty files
if(emptyFile)
{
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << "test, "
<< "atts, "
<< "units, "
<< "median, "
<< "mean, "
<< "stddev, "
<< "min, "
<< "max, ";
out << endl;
// Add headers only for empty files
if (emptyFile) {
// TODO: in big parallel runs, the "trials" are the procs
// and we really don't want to print them all out....
out << "test, "
<< "atts, "
<< "units, "
<< "median, "
<< "mean, "
<< "stddev, "
<< "min, "
<< "max, ";
out << endl;
}
for (int i=0; i<sorted.size(); i++)
{
Result &r = sorted[i];
for (int i = 0; i < sorted.size(); i++) {
Result& r = sorted[i];
out << r.test << ", ";
out << r.atts << ", ";
out << r.unit << ", ";
@@ -434,7 +374,7 @@ void ResultDatabase::DumpCsv(string fileName)
if (r.GetMean() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMean() << ", ";
out << r.GetMean() << ", ";
if (r.GetStdDev() == FLT_MAX)
out << "N/A, ";
else
@@ -442,11 +382,11 @@ void ResultDatabase::DumpCsv(string fileName)
if (r.GetMin() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMin() << ", ";
out << r.GetMin() << ", ";
if (r.GetMax() == FLT_MAX)
out << "N/A, ";
else
out << r.GetMax() << ", ";
out << r.GetMax() << ", ";
out << endl;
}
@@ -471,29 +411,24 @@ void ResultDatabase::DumpCsv(string fileName)
//
// ****************************************************************************
bool ResultDatabase::IsFileEmpty(string fileName)
{
bool fileEmpty;
bool ResultDatabase::IsFileEmpty(string fileName) {
bool fileEmpty;
ifstream file(fileName.c_str());
ifstream file(fileName.c_str());
//If the file doesn't exist it is by definition empty
if(!file.good())
{
// If the file doesn't exist it is by definition empty
if (!file.good()) {
return true;
}
else
{
} else {
fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
file.close();
return fileEmpty;
}
//Otherwise, return false
return false;
}
return fileEmpty;
}
// Otherwise, return false
return false;
}
// ****************************************************************************
@@ -511,16 +446,12 @@ bool ResultDatabase::IsFileEmpty(string fileName)
// Modifications:
//
// ****************************************************************************
vector<ResultDatabase::Result>
ResultDatabase::GetResultsForTest(const string &test)
{
vector<ResultDatabase::Result> ResultDatabase::GetResultsForTest(const string& test) {
// get only the given test results
vector<Result> retval;
for (int i=0; i<results.size(); i++)
{
Result &r = results[i];
if (r.test == test)
retval.push_back(r);
for (int i = 0; i < results.size(); i++) {
Result& r = results[i];
if (r.test == test) retval.push_back(r);
}
return retval;
}
@@ -539,8 +470,4 @@ ResultDatabase::GetResultsForTest(const string &test)
// Modifications:
//
// ****************************************************************************
const vector<ResultDatabase::Result> &
ResultDatabase::GetResults() const
{
return results;
}
const vector<ResultDatabase::Result>& ResultDatabase::GetResults() const { return results; }
projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.h
+22 -33
Zobrazit soubor
@@ -6,11 +6,11 @@
#include <iostream>
#include <fstream>
#include <cfloat>
using std::ifstream;
using std::ofstream;
using std::ostream;
using std::string;
using std::vector;
using std::ostream;
using std::ofstream;
using std::ifstream;
// ****************************************************************************
@@ -40,18 +40,16 @@ using std::ifstream;
// Added a GetResults method as well, and made several functions const.
//
// ****************************************************************************
class ResultDatabase
{
public:
class ResultDatabase {
public:
//
// A performance result for a single SHOC benchmark run.
//
struct Result
{
string test; // e.g. "readback"
string atts; // e.g. "pagelocked 4k^2"
string unit; // e.g. "MB/sec"
vector<double> value; // e.g. "837.14"
struct Result {
string test; // e.g. "readback"
string atts; // e.g. "pagelocked 4k^2"
string unit; // e.g. "MB/sec"
vector<double> value; // e.g. "837.14"
double GetMin() const;
double GetMax() const;
double GetMedian() const;
@@ -59,41 +57,32 @@ class ResultDatabase
double GetMean() const;
double GetStdDev() const;
bool operator<(const Result &rhs) const;
bool operator<(const Result& rhs) const;
bool HadAnyFLTMAXValues() const
{
for (int i=0; i<value.size(); ++i)
{
if (value[i] >= FLT_MAX)
return true;
bool HadAnyFLTMAXValues() const {
for (int i = 0; i < value.size(); ++i) {
if (value[i] >= FLT_MAX) return true;
}
return false;
}
};
protected:
protected:
vector<Result> results;
public:
void AddResult(const string &test,
const string &atts,
const string &unit,
double value);
void AddResults(const string &test,
const string &atts,
const string &unit,
const vector<double> &values);
vector<Result> GetResultsForTest(const string &test);
const vector<Result> &GetResults() const;
public:
void AddResult(const string& test, const string& atts, const string& unit, double value);
void AddResults(const string& test, const string& atts, const string& unit,
const vector<double>& values);
vector<Result> GetResultsForTest(const string& test);
const vector<Result>& GetResults() const;
void ClearAllResults();
void DumpDetailed(ostream&);
void DumpSummary(ostream&);
void DumpCsv(string fileName);
private:
private:
bool IsFileEmpty(string fileName);
};
projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+41 -46
Zobrazit soubor
@@ -21,35 +21,34 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
#include<iostream>
#include<time.h>
#include"ResultDatabase.h"
#include <iostream>
#include <time.h>
#include "ResultDatabase.h"
#define PRINT_PROGRESS 0
#define check(cmd) \
{\
hipError_t status = cmd;\
if(status != hipSuccess){ \
printf("error: '%s'(%d) from %s at %s:%d\n", \
hipGetErrorString(status), status, #cmd,\
__FILE__, __LINE__); \
abort(); \
}\
}
#define check(cmd) \
{ \
hipError_t status = cmd; \
if (status != hipSuccess) { \
printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(status), status, #cmd, \
__FILE__, __LINE__); \
abort(); \
} \
}
#define LEN 1024*1024
#define LEN 1024 * 1024
#define NUM_GROUPS 1
#define GROUP_SIZE 64
#define TEST_ITERS 20
#define TEST_ITERS 20
#define DISPATCHES_PER_TEST 100
const unsigned p_tests = 0xfffffff;
// HCC optimizes away fully NULL kernel calls, so run one that is nearly null:
__global__ void NearlyNull(hipLaunchParm lp, float* Ad){
__global__ void NearlyNull(hipLaunchParm lp, float* Ad) {
if (Ad) {
Ad[0] = 42;
}
@@ -59,38 +58,35 @@ __global__ void NearlyNull(hipLaunchParm lp, float* Ad){
ResultDatabase resultDB;
void stopTest(hipEvent_t start, hipEvent_t stop, const char *msg, int iters)
{
float mS = 0;
void stopTest(hipEvent_t start, hipEvent_t stop, const char* msg, int iters) {
float mS = 0;
check(hipEventRecord(stop));
check(hipDeviceSynchronize());
check(hipEventElapsedTime(&mS, start, stop));
resultDB.AddResult(std::string(msg), "", "uS", mS*1000/iters);
if (PRINT_PROGRESS & 0x1 ) {
std::cout<< msg <<"\t\t"<<mS*1000/iters<<" uS"<<std::endl;
resultDB.AddResult(std::string(msg), "", "uS", mS * 1000 / iters);
if (PRINT_PROGRESS & 0x1) {
std::cout << msg << "\t\t" << mS * 1000 / iters << " uS" << std::endl;
}
if (PRINT_PROGRESS & 0x2 ) {
if (PRINT_PROGRESS & 0x2) {
resultDB.DumpSummary(std::cout);
}
}
int main(){
hipError_t err;
float *Ad;
int main() {
hipError_t err;
float* Ad;
check(hipMalloc(&Ad, 4));
hipStream_t stream;
check(hipStreamCreate(&stream));
hipStream_t stream;
check(hipStreamCreate(&stream));
hipEvent_t start, sync, stop;
check(hipEventCreate(&start));
check(hipEventCreateWithFlags(&sync, hipEventBlockingSync));
check(hipEventCreate(&stop));
hipEvent_t start, sync, stop;
check(hipEventCreate(&start));
check(hipEventCreateWithFlags(&sync, hipEventBlockingSync));
check(hipEventCreate(&stop));
hipStream_t stream0 = 0;
@@ -103,7 +99,6 @@ int main(){
}
if (p_tests & 0x2) {
hipEventRecord(start);
hipLaunchKernel(NearlyNull, dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0, Ad);
@@ -112,9 +107,9 @@ int main(){
if (p_tests & 0x4) {
for (int t=0; t<TEST_ITERS; t++) {
for (int t = 0; t < TEST_ITERS; t++) {
hipEventRecord(start);
for(int i=0;i<DISPATCHES_PER_TEST;i++){
for (int i = 0; i < DISPATCHES_PER_TEST; i++) {
hipLaunchKernel(NearlyNull, dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0, Ad);
hipEventRecord(sync);
hipEventSynchronize(sync);
@@ -125,9 +120,9 @@ int main(){
if (p_tests & 0x10) {
for (int t=0; t<TEST_ITERS; t++) {
for (int t = 0; t < TEST_ITERS; t++) {
hipEventRecord(start);
for(int i=0;i<DISPATCHES_PER_TEST;i++){
for (int i = 0; i < DISPATCHES_PER_TEST; i++) {
hipLaunchKernel(NearlyNull, dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream, Ad);
hipEventRecord(sync);
hipEventSynchronize(sync);
@@ -139,9 +134,9 @@ int main(){
#if 1
if (p_tests & 0x40) {
for (int t=0; t<TEST_ITERS; t++) {
for (int t = 0; t < TEST_ITERS; t++) {
hipEventRecord(start);
for(int i=0;i<DISPATCHES_PER_TEST;i++){
for (int i = 0; i < DISPATCHES_PER_TEST; i++) {
hipLaunchKernel(NearlyNull, dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0, Ad);
}
stopTest(start, stop, "NullStreamASyncDispatchNoWait", DISPATCHES_PER_TEST);
@@ -149,9 +144,9 @@ int main(){
}
if (p_tests & 0x80) {
for (int t=0; t<TEST_ITERS; t++) {
for (int t = 0; t < TEST_ITERS; t++) {
hipEventRecord(start);
for(int i=0;i<DISPATCHES_PER_TEST;i++){
for (int i = 0; i < DISPATCHES_PER_TEST; i++) {
hipLaunchKernel(NearlyNull, dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream, Ad);
}
stopTest(start, stop, "StreamASyncDispatchNoWait", DISPATCHES_PER_TEST);
@@ -161,7 +156,7 @@ int main(){
resultDB.DumpSummary(std::cout);
check(hipEventDestroy(start));
check(hipEventDestroy(sync));
check(hipEventDestroy(stop));
check(hipEventDestroy(start));
check(hipEventDestroy(sync));
check(hipEventDestroy(stop));
}
projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp
+70 -68
Zobrazit soubor
@@ -24,61 +24,57 @@ THE SOFTWARE.
#include <iomanip>
#include "hip/hip_runtime.h"
#define KNRM "\x1B[0m"
#define KRED "\x1B[31m"
#define KGRN "\x1B[32m"
#define KYEL "\x1B[33m"
#define KBLU "\x1B[34m"
#define KMAG "\x1B[35m"
#define KCYN "\x1B[36m"
#define KWHT "\x1B[37m"
#define KNRM "\x1B[0m"
#define KRED "\x1B[31m"
#define KGRN "\x1B[32m"
#define KYEL "\x1B[33m"
#define KBLU "\x1B[34m"
#define KMAG "\x1B[35m"
#define KCYN "\x1B[36m"
#define KWHT "\x1B[37m"
#define failed(...) \
printf ("%serror: ", KRED);\
printf (__VA_ARGS__);\
printf ("\n");\
printf ("error: TEST FAILED\n%s", KNRM );\
#define failed(...) \
printf("%serror: ", KRED); \
printf(__VA_ARGS__); \
printf("\n"); \
printf("error: TEST FAILED\n%s", KNRM); \
exit(EXIT_FAILURE);
#define HIPCHECK(error) \
if (error != hipSuccess) { \
printf("%serror: '%s'(%d) at %s:%d%s\n", \
KRED, hipGetErrorString(error), error,\
__FILE__, __LINE__,KNRM);\
failed("API returned error code.");\
#define HIPCHECK(error) \
if (error != hipSuccess) { \
printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, hipGetErrorString(error), error, __FILE__, \
__LINE__, KNRM); \
failed("API returned error code."); \
}
void printCompilerInfo ()
{
void printCompilerInfo() {
#ifdef __HCC__
printf ("compiler: hcc version=%s, workweek (YYWWD) = %u\n", __hcc_version__, __hcc_workweek__);
printf("compiler: hcc version=%s, workweek (YYWWD) = %u\n", __hcc_version__, __hcc_workweek__);
#endif
#ifdef __NVCC__
printf ("compiler: nvcc\n");
printf("compiler: nvcc\n");
#endif
}
double bytesToGB(size_t s)
{
return (double)s / (1024.0*1024.0*1024.0);
}
double bytesToGB(size_t s) { return (double)s / (1024.0 * 1024.0 * 1024.0); }
#define printLimit(w1, limit, units) \
{\
size_t val;\
cudaDeviceGetLimit(&val, limit);\
std::cout << setw(w1) << #limit": " << val << " " << units << std::endl;\
}
#define printLimit(w1, limit, units) \
{ \
size_t val; \
cudaDeviceGetLimit(&val, limit); \
std::cout << setw(w1) << #limit ": " << val << " " << units << std::endl; \
}
void printDeviceProp (int deviceId)
{
void printDeviceProp(int deviceId) {
using namespace std;
const int w1 = 34;
cout << left;
cout << setw(w1) << "--------------------------------------------------------------------------------" << endl;
cout << setw(w1)
<< "--------------------------------------------------------------------------------"
<< endl;
cout << setw(w1) << "device#" << deviceId << endl;
hipDeviceProp_t props;
@@ -88,16 +84,22 @@ void printDeviceProp (int deviceId)
cout << setw(w1) << "pciBusID: " << props.pciBusID << endl;
cout << setw(w1) << "pciDeviceID: " << props.pciDeviceID << endl;
cout << setw(w1) << "multiProcessorCount: " << props.multiProcessorCount << endl;
cout << setw(w1) << "maxThreadsPerMultiProcessor: " << props.maxThreadsPerMultiProcessor << endl;
cout << setw(w1) << "maxThreadsPerMultiProcessor: " << props.maxThreadsPerMultiProcessor
<< endl;
cout << setw(w1) << "isMultiGpuBoard: " << props.isMultiGpuBoard << endl;
cout << setw(w1) << "clockRate: " << (float)props.clockRate / 1000.0 << " Mhz" << endl;
cout << setw(w1) << "memoryClockRate: " << (float)props.memoryClockRate / 1000.0 << " Mhz" << endl;
cout << setw(w1) << "memoryClockRate: " << (float)props.memoryClockRate / 1000.0 << " Mhz"
<< endl;
cout << setw(w1) << "memoryBusWidth: " << props.memoryBusWidth << endl;
cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate / 1000.0 << " Mhz" << endl;
cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2) << bytesToGB(props.totalGlobalMem) << " GB" << endl;
cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2) << bytesToGB(props.maxSharedMemoryPerMultiProcessor) << " GB" << endl;
cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate / 1000.0
<< " Mhz" << endl;
cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2)
<< bytesToGB(props.totalGlobalMem) << " GB" << endl;
cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2)
<< bytesToGB(props.maxSharedMemoryPerMultiProcessor) << " GB" << endl;
cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl;
cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB" << endl;
cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB"
<< endl;
cout << setw(w1) << "regsPerBlock: " << props.regsPerBlock << endl;
cout << setw(w1) << "warpSize: " << props.warpSize << endl;
cout << setw(w1) << "l2CacheSize: " << props.l2CacheSize << endl;
@@ -112,29 +114,31 @@ void printDeviceProp (int deviceId)
cout << setw(w1) << "major: " << props.major << endl;
cout << setw(w1) << "minor: " << props.minor << endl;
cout << setw(w1) << "concurrentKernels: " << props.concurrentKernels << endl;
cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl;
cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch << endl;
cout << setw(w1) << "arch.hasSharedInt32Atomics: " << props.arch.hasSharedInt32Atomics << endl;
cout << setw(w1) << "arch.hasSharedFloatAtomicExch: " << props.arch.hasSharedFloatAtomicExch << endl;
cout << setw(w1) << "arch.hasFloatAtomicAdd: " << props.arch.hasFloatAtomicAdd << endl;
cout << setw(w1) << "arch.hasGlobalInt64Atomics: " << props.arch.hasGlobalInt64Atomics << endl;
cout << setw(w1) << "arch.hasSharedInt64Atomics: " << props.arch.hasSharedInt64Atomics << endl;
cout << setw(w1) << "arch.hasDoubles: " << props.arch.hasDoubles << endl;
cout << setw(w1) << "arch.hasWarpVote: " << props.arch.hasWarpVote << endl;
cout << setw(w1) << "arch.hasWarpBallot: " << props.arch.hasWarpBallot << endl;
cout << setw(w1) << "arch.hasWarpShuffle: " << props.arch.hasWarpShuffle << endl;
cout << setw(w1) << "arch.hasFunnelShift: " << props.arch.hasFunnelShift << endl;
cout << setw(w1) << "arch.hasThreadFenceSystem: " << props.arch.hasThreadFenceSystem << endl;
cout << setw(w1) << "arch.hasSyncThreadsExt: " << props.arch.hasSyncThreadsExt << endl;
cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl;
cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl;
cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl;
cout << setw(w1) << "gcnArch: " << props.gcnArch << endl;
cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl;
cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch
<< endl;
cout << setw(w1) << "arch.hasSharedInt32Atomics: " << props.arch.hasSharedInt32Atomics << endl;
cout << setw(w1) << "arch.hasSharedFloatAtomicExch: " << props.arch.hasSharedFloatAtomicExch
<< endl;
cout << setw(w1) << "arch.hasFloatAtomicAdd: " << props.arch.hasFloatAtomicAdd << endl;
cout << setw(w1) << "arch.hasGlobalInt64Atomics: " << props.arch.hasGlobalInt64Atomics << endl;
cout << setw(w1) << "arch.hasSharedInt64Atomics: " << props.arch.hasSharedInt64Atomics << endl;
cout << setw(w1) << "arch.hasDoubles: " << props.arch.hasDoubles << endl;
cout << setw(w1) << "arch.hasWarpVote: " << props.arch.hasWarpVote << endl;
cout << setw(w1) << "arch.hasWarpBallot: " << props.arch.hasWarpBallot << endl;
cout << setw(w1) << "arch.hasWarpShuffle: " << props.arch.hasWarpShuffle << endl;
cout << setw(w1) << "arch.hasFunnelShift: " << props.arch.hasFunnelShift << endl;
cout << setw(w1) << "arch.hasThreadFenceSystem: " << props.arch.hasThreadFenceSystem << endl;
cout << setw(w1) << "arch.hasSyncThreadsExt: " << props.arch.hasSyncThreadsExt << endl;
cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl;
cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl;
cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl;
cout << setw(w1) << "gcnArch: " << props.gcnArch << endl;
int deviceCnt;
hipGetDeviceCount(&deviceCnt);
cout << setw(w1) << "peers: ";
for (int i=0; i<deviceCnt; i++) {
for (int i = 0; i < deviceCnt; i++) {
int isPeer;
hipDeviceCanAccessPeer(&isPeer, i, deviceId);
if (isPeer) {
@@ -143,7 +147,7 @@ void printDeviceProp (int deviceId)
}
cout << endl;
cout << setw(w1) << "non-peers: ";
for (int i=0; i<deviceCnt; i++) {
for (int i = 0; i < deviceCnt; i++) {
int isPeer;
hipDeviceCanAccessPeer(&isPeer, i, deviceId);
if (!isPeer) {
@@ -164,8 +168,6 @@ void printDeviceProp (int deviceId)
#endif
cout << endl;
@@ -174,11 +176,11 @@ void printDeviceProp (int deviceId)
cout << fixed << setprecision(2);
cout << setw(w1) << "memInfo.total: " << bytesToGB(total) << " GB" << endl;
cout << setw(w1) << "memInfo.free: " << bytesToGB(free) << " GB (" << setprecision(0) << (float)free/total * 100.0 << "%)" << endl;
cout << setw(w1) << "memInfo.free: " << bytesToGB(free) << " GB (" << setprecision(0)
<< (float)free / total * 100.0 << "%)" << endl;
}
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
using namespace std;
cout << endl;
@@ -189,7 +191,7 @@ int main(int argc, char *argv[])
HIPCHECK(hipGetDeviceCount(&deviceCnt));
for (int i=0; i< deviceCnt; i++) {
for (int i = 0; i < deviceCnt; i++) {
printDeviceProp(i);
}
projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp
+67 -80
Zobrazit soubor
@@ -20,28 +20,24 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define WIDTH 1024
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__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;
@@ -49,88 +45,79 @@ __global__ void matrixTranspose(hipLaunchParm lp,
}
// 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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
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 ) {
errors++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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);
// 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) {
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;
}
projects/hip-tests/samples/2_Cookbook/10_inline_asm/inline_asm.cpp
+106 -120
Zobrazit soubor
@@ -20,155 +20,141 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define WIDTH 1024
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__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]));
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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
float eventMs = 1.0f;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
float eventMs = 1.0f;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Record the start event
hipEventRecord(start, NULL);
return errors;
// 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;
}
projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
+4 -7
Zobrazit soubor
@@ -23,11 +23,8 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
extern texture<float, 2, hipReadModeElementType> tex;
__global__ void tex2dKernel(hipLaunchParm lp, float* outputData,
int width,
int height)
{
int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
outputData[y*width + x] = tex2D(tex, x, y);
__global__ void tex2dKernel(hipLaunchParm lp, float* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(tex, x, y);
}
projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+62 -60
Zobrazit soubor
@@ -32,111 +32,113 @@ THE SOFTWARE.
texture<float, 2, hipReadModeElementType> tex;
bool testResult = false;
#define HIP_CHECK(cmd) \
{\
hipError_t status = cmd;\
if(status != hipSuccess) {std::cout<<"error: #"<<status<<" ("<< hipGetErrorString(status) << ") at line:"<<__LINE__<<": "<<#cmd<<std::endl;abort();}\
}
#define HIP_CHECK(cmd) \
{ \
hipError_t status = cmd; \
if (status != hipSuccess) { \
std::cout << "error: #" << status << " (" << hipGetErrorString(status) \
<< ") at line:" << __LINE__ << ": " << #cmd << std::endl; \
abort(); \
} \
}
bool runTest(int argc, char **argv)
{
bool runTest(int argc, char** argv) {
unsigned int width = 256;
unsigned int height = 256;
unsigned int size = width * height * sizeof(float);
float* hData = (float*) malloc(size);
float* hData = (float*)malloc(size);
memset(hData, 0, size);
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
hData[i*width+j] = i*width+j;
hData[i * width + j] = i * width + j;
}
}
hipModule_t Module;
HIP_CHECK(hipModuleLoad(&Module, fileName));
hipArray* array;
HIP_ARRAY_DESCRIPTOR desc;
desc.format = HIP_AD_FORMAT_FLOAT;
desc.numChannels = 1;
desc.width = width;
desc.height = height;
HIP_ARRAY_DESCRIPTOR desc;
desc.format = HIP_AD_FORMAT_FLOAT;
desc.numChannels = 1;
desc.width = width;
desc.height = height;
hipArrayCreate(&array, &desc);
hip_Memcpy2D copyParam;
memset(&copyParam, 0, sizeof(copyParam));
copyParam.dstMemoryType = hipMemoryTypeArray;
copyParam.dstArray = array;
copyParam.srcMemoryType = hipMemoryTypeHost;
copyParam.srcHost = hData;
copyParam.srcPitch = width * sizeof(float);
copyParam.widthInBytes = copyParam.srcPitch;
copyParam.height = height;
memset(&copyParam, 0, sizeof(copyParam));
copyParam.dstMemoryType = hipMemoryTypeArray;
copyParam.dstArray = array;
copyParam.srcMemoryType = hipMemoryTypeHost;
copyParam.srcHost = hData;
copyParam.srcPitch = width * sizeof(float);
copyParam.widthInBytes = copyParam.srcPitch;
copyParam.height = height;
hipMemcpyParam2D(&copyParam);
textureReference* texref;
hipModuleGetTexRef(&texref, Module, "tex");
hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
hipTexRefSetFilterMode(texref, hipFilterModePoint);
hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
hipTexRefSetFilterMode(texref, hipFilterModePoint);
hipTexRefSetFlags(texref, 0);
hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
float* dData = NULL;
hipMalloc((void **) &dData, size);
hipMalloc((void**)&dData, size);
#ifdef __HIP_PLATFORM_HCC__
struct {
uint32_t _hidden[6]; // genco path + wrapper-gen pass used hidden arguments.
void * _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
struct {
uint32_t _hidden[6]; // genco path + wrapper-gen pass used hidden arguments.
void* _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
args._Ad = dData;
args._Bd = width;
args._Cd = height;
args._Bd = width;
args._Cd = height;
#endif
#ifdef __HIP_PLATFORM_NVCC__
struct {
uint32_t _hidden[1];
void * _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
struct {
uint32_t _hidden[1];
void* _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
args._hidden[0] = 0;
args._Ad = dData;
args._hidden[0] = 0;
args._Ad = dData;
args._Bd = width;
args._Cd = height;
args._Cd = height;
#endif
size_t sizeTemp = sizeof(args);
size_t sizeTemp = sizeof(args);
void *config[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &sizeTemp,
HIP_LAUNCH_PARAM_END
};
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
&sizeTemp, HIP_LAUNCH_PARAM_END};
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel"));
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel"));
int temp1= width/16;
int temp2 = height/16;
HIP_CHECK(hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config));
int temp1 = width / 16;
int temp2 = height / 16;
HIP_CHECK(
hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config));
hipDeviceSynchronize();
float *hOutputData = (float *) malloc(size);
memset(hOutputData, 0, size);
float* hOutputData = (float*)malloc(size);
memset(hOutputData, 0, size);
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
if (hData[i*width+j] != hOutputData[i*width+j]) {
printf("Difference [ %d %d ]:%f ----%f\n",i, j, hData[i*width+j] , hOutputData[i*width+j]);
if (hData[i * width + j] != hOutputData[i * width + j]) {
printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j],
hOutputData[i * width + j]);
testResult = false;
break;
}
@@ -147,7 +149,7 @@ bool runTest(int argc, char **argv)
return true;
}
int main(int argc, char **argv){
int main(int argc, char** argv) {
hipInit(0);
testResult = runTest(argc, argv);
printf("%s ...\n", testResult ? "PASSED" : "FAILED");
projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp
+67 -80
Zobrazit soubor
@@ -20,28 +20,24 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define WIDTH 1024
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__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;
@@ -49,88 +45,79 @@ __global__ void matrixTranspose(hipLaunchParm lp,
}
// 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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
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 ) {
errors++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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);
// 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) {
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;
}
projects/hip-tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp
+104 -117
Zobrazit soubor
@@ -20,26 +20,22 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define WIDTH 1024
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__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;
@@ -47,126 +43,117 @@ __global__ void matrixTranspose(hipLaunchParm lp,
}
// 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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
float eventMs = 1.0f;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
float eventMs = 1.0f;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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 ) {
errors++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Record the start event
hipEventRecord(start, NULL);
return errors;
// 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) {
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;
}
projects/hip-tests/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
+159 -172
Zobrazit soubor
@@ -20,33 +20,29 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#include "hip/hip_profile.h"
#define WIDTH 1024
#define WIDTH 1024
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#define ITERATIONS 10
// Cmdline parms to control start and stop triggers
int startTriggerIteration=-1;
int stopTriggerIteration=-1;
int startTriggerIteration = -1;
int stopTriggerIteration = -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)
{
__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;
@@ -54,180 +50,171 @@ __global__ void matrixTranspose(hipLaunchParm lp,
}
// 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];
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];
}
}
}
// Use a separate function to demonstrate how to use function name as part of scoped marker:
void runGPU(float *Matrix, float *TransposeMatrix,
float* gpuMatrix, float* gpuTransposeMatrix) {
void runGPU(float* Matrix, float* TransposeMatrix, float* gpuMatrix, float* gpuTransposeMatrix) {
// __func__ is a standard C++ macro which expands to the name of the function, in this case
// "runGPU"
HIP_SCOPED_MARKER(__func__, "MyGroup");
// __func__ is a standard C++ macro which expands to the name of the function, in this case "runGPU"
HIP_SCOPED_MARKER(__func__, "MyGroup");
for (int i = 0; i < ITERATIONS; i++) {
if (i == startTriggerIteration) {
hipProfilerStart();
}
if (i == stopTriggerIteration) {
hipProfilerStop();
}
for (int i=0; i<ITERATIONS; i++) {
float eventMs = 0.0f;
if (i==startTriggerIteration) {
hipProfilerStart();
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
// 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);
}
if (i==stopTriggerIteration) {
hipProfilerStop();
}
float eventMs = 0.0f;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
// 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);
}
};
int main(int argc, char *argv[]) {
if (argc >= 2) {
startTriggerIteration = atoi(argv[1]);
printf ("info : will start tracing at iteration:%d\n", startTriggerIteration);
}
if (argc >= 3) {
stopTriggerIteration = atoi(argv[2]);
printf ("info : will stop tracing at iteration:%d\n", stopTriggerIteration);
}
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
{
// Show example of how to create a "scoped marker".
// The scoped marker records the time spent inside the { scope } of the marker - the begin timestamp is at the
// beginning of the code scope, and the end is recorded when the SCOPE exits. This can be viewed in CodeXL
// timeline relative to other GPU and CPU events.
// This marker captures the time spent in setup including host allocation, initialization, and device memory allocation.
HIP_SCOPED_MARKER("Setup", "MyGroup");
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
// initialize the input data
for (int 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));
// FYI, the scoped-marker will be destroyed here when the scope exits, and will record its "end" timestamp.
}
runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix);
// show how to use explicit begin/end markers:
// We begin the timed region with HIP_BEGIN_MARKER, passing in the markerName and group:
// The region will stop when HIP_END_MARKER is called
// This is another way to mark begin/end - as an alternative to scoped markers.
HIP_BEGIN_MARKER("Check&TearDown", "MyGroup");
int errors = 0;
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
// verify the results
double eps = 1.0E-6;
for (int i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
errors++;
int main(int argc, char* argv[]) {
if (argc >= 2) {
startTriggerIteration = atoi(argv[1]);
printf("info : will start tracing at iteration:%d\n", startTriggerIteration);
}
if (argc >= 3) {
stopTriggerIteration = atoi(argv[2]);
printf("info : will stop tracing at iteration:%d\n", stopTriggerIteration);
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
float* gpuMatrix;
float* gpuTransposeMatrix;
// This ends the last marker started in this thread, in this case "Check&TearDown"
HIP_END_MARKER();
return errors;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
{
// Show example of how to create a "scoped marker".
// The scoped marker records the time spent inside the { scope } of the marker - the begin
// timestamp is at the beginning of the code scope, and the end is recorded when the SCOPE
// exits. This can be viewed in CodeXL timeline relative to other GPU and CPU events. This
// marker captures the time spent in setup including host allocation, initialization, and
// device memory allocation.
HIP_SCOPED_MARKER("Setup", "MyGroup");
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
// initialize the input data
for (int 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));
// FYI, the scoped-marker will be destroyed here when the scope exits, and will record its
// "end" timestamp.
}
runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix);
// show how to use explicit begin/end markers:
// We begin the timed region with HIP_BEGIN_MARKER, passing in the markerName and group:
// The region will stop when HIP_END_MARKER is called
// This is another way to mark begin/end - as an alternative to scoped markers.
HIP_BEGIN_MARKER("Check&TearDown", "MyGroup");
int errors = 0;
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
// verify the results
double eps = 1.0E-6;
for (int i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
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);
// This ends the last marker started in this thread, in this case "Check&TearDown"
HIP_END_MARKER();
return errors;
}
projects/hip-tests/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp
+69 -82
Zobrazit soubor
@@ -20,28 +20,24 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 64
#define WIDTH 64
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__shared__ float sharedMem[WIDTH*WIDTH];
__global__ void matrixTranspose(hipLaunchParm lp, float* out, float* in, const int width) {
__shared__ float sharedMem[WIDTH * WIDTH];
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
@@ -54,89 +50,80 @@ __global__ void matrixTranspose(hipLaunchParm lp,
}
// 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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
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++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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);
// 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;
}
projects/hip-tests/samples/2_Cookbook/4_shfl/shfl.cpp
+69 -85
Zobrazit soubor
@@ -20,122 +20,106 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 4
#define WIDTH 4
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__global__ void matrixTranspose(hipLaunchParm lp, float* out, float* in, const int width) {
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
float val = in[x];
for(int i=0;i<width;i++)
{
for(int j=0;j<width;j++)
out[i*width + j] = __shfl(val,j*width + i);
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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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;
}
projects/hip-tests/samples/2_Cookbook/5_2dshfl/2dshfl.cpp
+69 -83
Zobrazit soubor
@@ -20,118 +20,104 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 4
#define WIDTH 4
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__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;
float val = in[y*width + x];
float val = in[y * width + x];
out[x*width + y] = __shfl(val,y*width + x);
out[x * width + y] = __shfl(val, 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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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;
}
projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp
+68 -81
Zobrazit soubor
@@ -19,26 +19,22 @@ 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>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 16
#define WIDTH 16
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__global__ void matrixTranspose(hipLaunchParm lp, float* out, float* in, const int width) {
// declare dynamic shared memory
HIP_DYNAMIC_SHARED(float, sharedMem);
@@ -53,89 +49,80 @@ __global__ void matrixTranspose(hipLaunchParm lp,
}
// 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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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(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);
// 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++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("dynamic_shared PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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), sizeof(float) * WIDTH * WIDTH,
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("dynamic_shared 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;
}
projects/hip-tests/samples/2_Cookbook/7_streams/stream.cpp
+36 -48
Zobrazit soubor
@@ -20,22 +20,19 @@ THE SOFTWARE.
#include <iostream>
#include <hip/hip_runtime.h>
#define WIDTH 32
#define WIDTH 32
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
using namespace std;
__global__ void matrixTranspose_static_shared(hipLaunchParm lp,
float *out,
float *in,
const int width)
{
__shared__ float sharedMem[WIDTH*WIDTH];
__global__ void matrixTranspose_static_shared(hipLaunchParm lp, float* out, float* in,
const int width) {
__shared__ float sharedMem[WIDTH * WIDTH];
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
@@ -47,11 +44,8 @@ __global__ void matrixTranspose_static_shared(hipLaunchParm lp,
out[y * width + x] = sharedMem[y * width + x];
}
__global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
float *out,
float *in,
const int width)
{
__global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp, float* out, float* in,
const int width) {
// declare dynamic shared memory
HIP_DYNAMIC_SHARED(float, sharedMem)
@@ -65,39 +59,34 @@ __global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
out[y * width + x] = sharedMem[y * width + x];
}
void MultipleStream (float **data, float *randArray, float **gpuTransposeMatrix, float **TransposeMatrix, int width)
{
void MultipleStream(float** data, float* randArray, float** gpuTransposeMatrix,
float** TransposeMatrix, int width) {
const int num_streams = 2;
hipStream_t streams[num_streams];
for(int i=0;i<num_streams;i++)
hipStreamCreate(&streams[i]);
for (int i = 0; i < num_streams; i++) hipStreamCreate(&streams[i]);
for(int i=0;i<num_streams;i++)
{
for (int i = 0; i < num_streams; i++) {
hipMalloc((void**)&data[i], NUM * sizeof(float));
hipMemcpyAsync(data[i], randArray, NUM * sizeof(float), hipMemcpyHostToDevice,streams[i]);
hipMemcpyAsync(data[i], randArray, NUM * sizeof(float), hipMemcpyHostToDevice, streams[i]);
}
hipLaunchKernel(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],
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);
hipLaunchKernel(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);
for(int i=0;i<num_streams;i++)
hipMemcpyAsync(TransposeMatrix[i], gpuTransposeMatrix[i], NUM*sizeof(float), hipMemcpyDeviceToHost, streams[i]);
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);
for (int i = 0; i < num_streams; i++)
hipMemcpyAsync(TransposeMatrix[i], gpuTransposeMatrix[i], NUM * sizeof(float),
hipMemcpyDeviceToHost, streams[i]);
}
int main(){
int main() {
hipSetDevice(0);
float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray;
@@ -112,9 +101,8 @@ int main(){
hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
for(int i = 0; i < NUM; i++)
{
randArray[i] = (float)i*1.0f;
for (int i = 0; i < NUM; i++) {
randArray[i] = (float)i * 1.0f;
}
MultipleStream(data, randArray, gpuTransposeMatrix, TransposeMatrix, width);
@@ -125,22 +113,22 @@ int main(){
int errors = 0;
double eps = 1.0E-6;
for (int i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[0][i] - TransposeMatrix[1][i]) > eps ) {
printf("%d stream0: %f stream1 %f\n",i,TransposeMatrix[0][i],TransposeMatrix[1][i]);
errors++;
if (std::abs(TransposeMatrix[0][i] - TransposeMatrix[1][i]) > eps) {
printf("%d stream0: %f stream1 %f\n", i, TransposeMatrix[0][i], TransposeMatrix[1][i]);
errors++;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
if (errors != 0) {
printf("FAILED: %d errors\n", errors);
} else {
printf ("stream PASSED!\n");
printf("stream PASSED!\n");
}
free(randArray);
for(int i=0;i<2;i++){
hipFree(data[i]);
hipFree(gpuTransposeMatrix[i]);
free(TransposeMatrix[i]);
for (int i = 0; i < 2; i++) {
hipFree(data[i]);
hipFree(gpuTransposeMatrix[i]);
free(TransposeMatrix[i]);
}
hipDeviceReset();
projects/hip-tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp
+67 -85
Zobrazit soubor
@@ -20,105 +20,94 @@ THE SOFTWARE.
#include <iostream>
#include <hip/hip_runtime.h>
#include <assert.h>
#define WIDTH 32
#define WIDTH 32
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
using namespace std;
#define KNRM "\x1B[0m"
#define KRED "\x1B[31m"
#define KNRM "\x1B[0m"
#define KRED "\x1B[31m"
#define failed(...) \
printf ("%serror: ", KRED);\
printf (__VA_ARGS__);\
printf ("\n");\
printf ("error: TEST FAILED\n%s", KNRM );\
#define failed(...) \
printf("%serror: ", KRED); \
printf(__VA_ARGS__); \
printf("\n"); \
printf("error: TEST FAILED\n%s", KNRM); \
abort();
#define HIPCHECK(error) \
{\
hipError_t localError = error; \
if (localError != hipSuccess) { \
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
KRED, hipGetErrorString(localError), localError,\
#error,__FILE__, __LINE__, KNRM); \
failed("API returned error code.");\
}\
}
#define HIPCHECK(error) \
{ \
hipError_t localError = error; \
if (localError != hipSuccess) { \
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError), \
localError, #error, __FILE__, __LINE__, KNRM); \
failed("API returned error code."); \
} \
}
void checkPeer2PeerSupport()
{
void checkPeer2PeerSupport() {
int gpuCount;
int canAccessPeer;
HIPCHECK(hipGetDeviceCount(&gpuCount));
for (int currentGpu=0; currentGpu<gpuCount; currentGpu++)
{
for (int currentGpu = 0; currentGpu < gpuCount; currentGpu++) {
HIPCHECK(hipSetDevice(currentGpu));
for (int peerGpu=0; peerGpu<currentGpu; peerGpu++)
{
if (currentGpu!=peerGpu)
{
for (int peerGpu = 0; peerGpu < currentGpu; peerGpu++) {
if (currentGpu != peerGpu) {
HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu));
printf ("currentGpu#%d canAccessPeer: peerGpu#%d=%d\n", currentGpu, peerGpu, canAccessPeer);
printf("currentGpu#%d canAccessPeer: peerGpu#%d=%d\n", currentGpu, peerGpu,
canAccessPeer);
}
HIPCHECK(hipSetDevice(peerGpu));
HIPCHECK(hipDeviceReset());
}
HIPCHECK(hipSetDevice(currentGpu));
HIPCHECK(hipDeviceReset());
HIPCHECK(hipSetDevice(currentGpu));
HIPCHECK(hipDeviceReset());
}
}
void enablePeer2Peer(int currentGpu, int peerGpu)
{
void enablePeer2Peer(int currentGpu, int peerGpu) {
int canAccessPeer;
// Must be on a multi-gpu system:
assert (currentGpu != peerGpu);
assert(currentGpu != peerGpu);
HIPCHECK(hipSetDevice(currentGpu));
hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
if(canAccessPeer==1){
if (canAccessPeer == 1) {
HIPCHECK(hipDeviceEnablePeerAccess(peerGpu, 0));
}
else
printf("peer2peer transfer not possible between the selected gpu devices");
} else
printf("peer2peer transfer not possible between the selected gpu devices");
}
void disablePeer2Peer(int currentGpu, int peerGpu)
{
void disablePeer2Peer(int currentGpu, int peerGpu) {
int canAccessPeer;
// Must be on a multi-gpu system:
assert (currentGpu != peerGpu);
assert(currentGpu != peerGpu);
HIPCHECK(hipSetDevice(currentGpu));
hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
if(canAccessPeer==1){
if (canAccessPeer == 1) {
HIPCHECK(hipDeviceDisablePeerAccess(peerGpu));
}
else
printf("peer2peer disable not required");
} else
printf("peer2peer disable not required");
}
__global__ void matrixTranspose_static_shared(hipLaunchParm lp,
float *out,
float *in,
const int width)
{
__shared__ float sharedMem[WIDTH*WIDTH];
__global__ void matrixTranspose_static_shared(hipLaunchParm lp, float* out, float* in,
const int width) {
__shared__ float sharedMem[WIDTH * WIDTH];
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
@@ -130,11 +119,8 @@ __global__ void matrixTranspose_static_shared(hipLaunchParm lp,
out[y * width + x] = sharedMem[y * width + x];
}
__global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
float *out,
float *in,
const int width)
{
__global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp, float* out, float* in,
const int width) {
// declare dynamic shared memory
HIP_DYNAMIC_SHARED(float, sharedMem)
@@ -148,8 +134,7 @@ __global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
out[y * width + x] = sharedMem[y * width + x];
}
int main(){
int main() {
checkPeer2PeerSupport();
int gpuCount;
@@ -157,8 +142,7 @@ int main(){
HIPCHECK(hipGetDeviceCount(&gpuCount));
if (gpuCount < 2)
{
if (gpuCount < 2) {
printf("Peer2Peer application requires atleast 2 gpu devices");
return 0;
}
@@ -166,7 +150,7 @@ int main(){
currentGpu = 0;
peerGpu = (currentGpu + 1);
printf ("currentGpu=%d peerGpu=%d (Total no. of gpu = %d)\n", currentGpu, peerGpu, gpuCount);
printf("currentGpu=%d peerGpu=%d (Total no. of gpu = %d)\n", currentGpu, peerGpu, gpuCount);
float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray;
@@ -174,9 +158,8 @@ int main(){
randArray = (float*)malloc(NUM * sizeof(float));
for(int i = 0; i < NUM; i++)
{
randArray[i] = (float)i*1.0f;
for (int i = 0; i < NUM; i++) {
randArray[i] = (float)i * 1.0f;
}
enablePeer2Peer(currentGpu, peerGpu);
@@ -188,10 +171,9 @@ int main(){
hipMemcpy(data[0], randArray, NUM * sizeof(float), hipMemcpyHostToDevice);
hipLaunchKernel(matrixTranspose_static_shared,
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix[0], data[0], width);
dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix[0],
data[0], width);
HIPCHECK(hipSetDevice(peerGpu));
TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
@@ -200,12 +182,12 @@ int main(){
hipMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), hipMemcpyDeviceToDevice);
hipLaunchKernel(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, 0,
gpuTransposeMatrix[1], data[1], width);
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[1], data[1], width);
hipMemcpy(TransposeMatrix[1], gpuTransposeMatrix[1], NUM*sizeof(float), hipMemcpyDeviceToHost);
hipMemcpy(TransposeMatrix[1], gpuTransposeMatrix[1], NUM * sizeof(float),
hipMemcpyDeviceToHost);
hipDeviceSynchronize();
@@ -215,22 +197,22 @@ int main(){
int errors = 0;
double eps = 1.0E-6;
for (int i = 0; i < NUM; i++) {
if (std::abs(randArray[i] - TransposeMatrix[1][i]) > eps ) {
printf("%d cpu: %f gpu peered data %f\n",i,randArray[i],TransposeMatrix[1][i]);
errors++;
if (std::abs(randArray[i] - TransposeMatrix[1][i]) > eps) {
printf("%d cpu: %f gpu peered data %f\n", i, randArray[i], TransposeMatrix[1][i]);
errors++;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
if (errors != 0) {
printf("FAILED: %d errors\n", errors);
} else {
printf ("Peer2Peer PASSED!\n");
printf("Peer2Peer PASSED!\n");
}
free(randArray);
for(int i=0;i<2;i++){
hipFree(data[i]);
hipFree(gpuTransposeMatrix[i]);
free(TransposeMatrix[i]);
for (int i = 0; i < 2; i++) {
hipFree(data[i]);
hipFree(gpuTransposeMatrix[i]);
free(TransposeMatrix[i]);
}
HIPCHECK(hipSetDevice(peerGpu));
projects/hip-tests/samples/2_Cookbook/9_unroll/unroll.cpp
+69 -85
Zobrazit soubor
@@ -20,122 +20,106 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<iostream>
#include <iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 4
#define WIDTH 4
#define NUM (WIDTH*WIDTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
#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)
{
__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);
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];
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* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
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++;
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
return errors;
// 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;
}