AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Apply .clangformat to all repo source files

Browse Source 
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344
This commit is contained in:
Maneesh Gupta
2018-03-12 11:29:03 +05:30
parent eee7fa6072
commit 9e47fccc89
293 changed files with 43980 additions and 45830 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/samples/0_Intro/bit_extract/bit_extract.cpp
+32 -35
View File
@@ -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");
}
hipamd/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp
+19 -18
View File
@@ -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");
}
hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
+26 -26
View File
@@ -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");
}
hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
+17 -18
View File
@@ -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");
}
hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
+15 -15
View File
@@ -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");
}
hipamd/samples/0_Intro/hcc_dialects/vadd_hip.cpp
+13 -15
View File
@@ -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");
}
hipamd/samples/0_Intro/module_api/defaultDriver.cpp
+10 -10
View File
@@ -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;
}
}
hipamd/samples/0_Intro/module_api/launchKernelHcc.cpp
+19 -20
View File
@@ -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;
}
}
hipamd/samples/0_Intro/module_api/runKernel.cpp
+19 -19
View File
@@ -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;
}
}
hipamd/samples/0_Intro/module_api/vcpy_kernel.cpp
+1 -2
View File
@@ -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];
}
hipamd/samples/0_Intro/module_api_global/runKernel.cpp
+30 -29
View File
@@ -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;
}
hipamd/samples/0_Intro/module_api_global/vcpy_kernel.cpp
+5 -7
View File
@@ -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];
}
hipamd/samples/0_Intro/square/square.hipref.cpp
+32 -35
View File
@@ -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");
}