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

Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

Browse Source 
Conflicts:
	src/hip_hcc.cpp
This commit is contained in:
Ben Sander
2016-03-19 03:22:09 -05:00
parent efc9df8805 a17e390283
commit 90ad8ddc5d
15 changed files with 605 additions and 60 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/tests/src/cudaRegister.cu
+90
View File
@@ -0,0 +1,90 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<cuda.h>
#include<cuda_runtime.h>
#include<iostream>
#include<unistd.h>
#include<stdio.h>
#include<malloc.h>
#define LEN 1024
#define SIZE LEN * sizeof(float)
#define ITER 1024*1024
#define check(msg, status){ \
if(status != cudaSuccess) { \
printf("%s failed. \n", #msg); \
} \
}
__global__ void Inc1(float *Ad, float *Bd){
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if(tx < 1 ){
for(int i=0;i<ITER;i++){
Ad[tx] = Ad[tx] + 1.0f;
for(int j=0;j<256;j++){
Bd[tx] = Ad[tx];
}
}
}
}
__global__ void Inc2(float *Ad, float *Bd){
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if(tx < 1024){
for(int i=0;i<ITER;i++){
Ad[tx] = Ad[tx] + 1.0f;
for(int j=0;j<256;j++){
Bd[tx] = Ad[tx];
}
}
}
}
int main(){
float *A, *Ad, *Bd;
A = new float[LEN];
for(int i=0;i<LEN;i++){
A[i] = 0.0f;
}
cudaError_t status;
status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped);
check("Registering A",status);
cudaHostGetDevicePointer(&Ad, A, 0);
cudaMalloc((void**)&Bd, SIZE);
dim3 dimGrid(LEN/512,1,1);
dim3 dimBlock(512,1,1);
Inc1<<<dimGrid, dimBlock>>>(Ad, Bd);
sleep(3);
A[0] = -(ITER*1.0f);
std::cout<<"Same cache line before completion: \t"<< A[0]<<std::endl;
cudaDeviceSynchronize();
std::cout<<"Same cache line after completion: \t"<< A[0]<<std::endl;
for(int i=0;i<LEN;i++){
A[i] = 0.0f;
}
Inc2<<<dimGrid, dimBlock>>>(Ad, Bd);
sleep(3);
A[0] = -(ITER*1.0f);
std::cout<<"Diff cache line before completion: \t"<<A[0]<<std::endl;
cudaDeviceSynchronize();
std::cout<<"Diff cache line after completion: \t"<<A[0]<<std::endl;
}
hipamd/tests/src/hipHostAlloc.cpp
+6 -3
View File
@@ -39,14 +39,14 @@ int device;
HIPCHECK(hipGetDevice(&device));
HIPCHECK(hipGetDeviceProperties(&prop, device));
if(prop.canMapHostMemory != 1){
std::cout<<"Exiting..."<<std::endl;
//std::cout<<"Exiting..."<<std::endl;
}
HIPCHECK(hipHostAlloc((void**)&A, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&C, SIZE, hipHostAllocMapped));
HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Bd, B, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
for(int i=0;i<LEN;i++){
@@ -54,6 +54,9 @@ A[i] = 1.0f;
B[i] = 2.0f;
}
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
dim3 dimGrid(LEN/512,1,1);
dim3 dimBlock(512,1,1);
hipamd/tests/src/hipMemcpy.cpp
+3 -3
View File
@@ -36,7 +36,7 @@ void printSep()
// The subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result.
//
// IN: numElements controls the number of elements used for allocations.
// IN: usePinnedHost : If true, allocate host with hipMallocHost and is pinned ; else allocate host memory with malloc.
// IN: usePinnedHost : If true, allocate host with hipHostAlloc and is pinned ; else allocate host memory with malloc.
// IN: useHostToHost : If true, add an extra host-to-host copy.
// IN: useDeviceToDevice : If true, add an extra deviceto-device copy after result is produced.
// IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction.
@@ -67,8 +67,8 @@ void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, boo
if (useHostToHost) {
if (usePinnedHost) {
HIPCHECK ( hipMallocHost(&A_hh, sizeElements) );
HIPCHECK ( hipMallocHost(&B_hh, sizeElements) );
HIPCHECK ( hipHostAlloc((void**)&A_hh, sizeElements, hipHostAllocDefault) );
HIPCHECK ( hipHostAlloc((void**)&B_hh, sizeElements, hipHostAllocDefault) );
} else {
A_hh = (T*)malloc(sizeElements);
B_hh = (T*)malloc(sizeElements);
hipamd/tests/src/hipMemcpyAsync.cpp
+4 -4
View File
@@ -33,7 +33,7 @@ void simpleNegTest()
size_t Nbytes = N*sizeof(float);
A_malloc = (float*)malloc(Nbytes);
HIPCHECK(hipMallocHost(&A_pinned, Nbytes));
HIPCHECK(hipHostAlloc((void**)&A_pinned, Nbytes, hipHostAllocDefault));
HIPCHECK(hipMalloc(&A_d, Nbytes));
@@ -61,7 +61,7 @@ struct HostTraits<Pinned>
static void *Alloc(size_t sizeBytes) {
void *p;
HIPCHECK(hipMallocHost(&p, sizeBytes));
HIPCHECK(hipHostAlloc((void**)&p, sizeBytes, hipHostAllocDefault));
return p;
};
};
@@ -181,8 +181,8 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
T *A_d;
T *A_h1, *A_h2;
HIPCHECK(hipMallocHost(&A_h1, Nbytes));
HIPCHECK(hipMallocHost(&A_h2, Nbytes));
HIPCHECK(hipHostAlloc((void**)&A_h1, Nbytes, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&A_h2, Nbytes, hipHostAllocDefault));
HIPCHECK(hipMalloc(&A_d, Nbytes));
for (int i=0; i<numElements; i++) {
hipamd/tests/src/hipMemcpy_simple.cpp
+2 -2
View File
@@ -68,8 +68,8 @@ void simpleTest2(size_t numElements, bool usePinnedHost)
T *A_d, *A_h1, *A_h2;
if (usePinnedHost) {
HIPCHECK ( hipMallocHost(&A_h1, sizeElements) );
HIPCHECK ( hipMallocHost(&A_h2, sizeElements) );
HIPCHECK ( hipHostAlloc((void**)&A_h1, sizeElements, hipHostAllocDefault) );
HIPCHECK ( hipHostAlloc((void**)&A_h2, sizeElements, hipHostAllocDefault) );
} else {
A_h1 = (T*)aligned_alloc(alignment, sizeElements);
HIPASSERT(A_h1);
hipamd/tests/src/hipMultiThreadStreams2.cpp
+27 -27
View File
@@ -34,21 +34,21 @@ Array[tx] = Array[tx] + T(1);
void run1(size_t size, hipStream_t stream){
float *Ah, *Bh, *Cd, *Dd, *Eh;
hipMallocHost(&Ah, size);
hipMallocHost(&Bh, size);
hipMalloc(&Cd, size);
hipMalloc(&Dd, size);
hipMallocHost(&Eh, size);
HIPCHECK(hipHostAlloc((void**)&Ah, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Bh, size, hipHostAllocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostAlloc((void**)&Eh, size, hipHostAllocDefault));
for(int i=0;i<N;i++){
Ah[i] = 1.0f;
}
hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream);
hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream);
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream));
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream);
hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream));
HIPCHECK(hipDeviceSynchronize());
HIPASSERT(Eh[10] == Ah[10] + 1.0f);
}
@@ -58,27 +58,27 @@ void run(size_t size, hipStream_t stream1, hipStream_t stream2){
float *Ah, *Bh, *Cd, *Dd, *Eh;
float *Ahh, *Bhh, *Cdd, *Ddd, *Ehh;
hipMallocHost(&Ah, size);
hipMallocHost(&Bh, size);
hipMalloc(&Cd, size);
hipMalloc(&Dd, size);
hipMallocHost(&Eh, size);
hipMallocHost(&Ahh, size);
hipMallocHost(&Bhh, size);
hipMalloc(&Cdd, size);
hipMalloc(&Ddd, size);
hipMallocHost(&Ehh, size);
HIPCHECK(hipHostAlloc((void**)&Ah, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Bh, size, hipHostAllocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostAlloc((void**)&Eh, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Ahh, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Bhh, size, hipHostAllocDefault));
HIPCHECK(hipMalloc(&Cdd, size));
HIPCHECK(hipMalloc(&Ddd, size));
HIPCHECK(hipHostAlloc((void**)&Ehh, size, hipHostAllocDefault));
hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1);
hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2);
hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1);
hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2);
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1));
HIPCHECK(hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2));
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream1, Cd);
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream2, Cdd);
hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1);
hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2);
hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1);
hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1));
HIPCHECK(hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2));
HIPCHECK(hipDeviceSynchronize());
HIPASSERT(Eh[10] = Ah[10] + 1.0f);
HIPASSERT(Ehh[10] = Ahh[10] + 1.0f);
hipamd/tests/src/hipPerfMemcpy.cpp
+41
View File
@@ -0,0 +1,41 @@
#include"test_common.h"
#include<iostream>
#include<time.h>
#define NUM_SIZE 8
#define NUM_ITER 12
static size_t size[NUM_SIZE];
void setup(){
for(int i=0;i<NUM_SIZE;i++){
size[i] = 1<<(i+6); // start at 8 bytes
}
}
void valSet(int *A, int val, size_t size){
size_t len = size/sizeof(int);
for(int i=0;i<len;i++){
A[i] = val;
}
}
int main(){
setup();
int *A, *Ad;
for(int i=0;i<NUM_SIZE;i++){
std::cout<<size[i]<<std::endl;
A = (int*)malloc(size[i]);
valSet(A, 1, size[i]);
hipMalloc(&Ad, size[i]);
std::cout<<"Malloc success at size: "<<size[i]<<std::endl;
clock_t start ,end;
start = clock();
for(int i=0;i<NUM_ITER;i++){
hipMemcpy(Ad, A, size[i], hipMemcpyHostToDevice);
}
hipDeviceSynchronize();
end = clock();
double uS = (double)(end - start)*1000/(NUM_ITER*CLOCKS_PER_SEC);
std::cout<<uS<<std::endl;
}
}
hipamd/tests/src/hipPointerAttrib.cpp
+4 -4
View File
@@ -115,7 +115,7 @@ void testSimple()
hipError_t e;
HIPCHECK ( hipMalloc(&A_d, Nbytes) );
HIPCHECK ( hipMallocHost(&A_Pinned_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)&A_Pinned_h, Nbytes, hipHostAllocDefault) );
A_OSAlloc_h = (char*)malloc(Nbytes);
size_t free, total;
@@ -168,7 +168,7 @@ void testSimple()
// Device-visible host memory
printf ("\nDevice-visible host memory (hipMallocHost)\n");
printf ("\nDevice-visible host memory (hipHostAlloc)\n");
HIPCHECK( hipPointerGetAttributes(&attribs, A_Pinned_h));
printf("getAttr:%-20s", "A_pinned_h"); printAttribs(&attribs);
@@ -277,13 +277,13 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
void * ptr;
if (isDevice) {
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
HIPCHECK(hipMalloc(&ptr, reference[i]._sizeBytes));
HIPCHECK(hipHostAlloc((void**)&ptr, reference[i]._sizeBytes, hipHostAllocDefault));
reference[i]._attrib.memoryType = hipMemoryTypeDevice;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = NULL;
reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
} else {
HIPCHECK(hipMallocHost(&ptr, reference[i]._sizeBytes));
HIPCHECK(hipHostAlloc((void**)&ptr, reference[i]._sizeBytes, hipHostAllocDefault));
reference[i]._attrib.memoryType = hipMemoryTypeHost;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = ptr;
hipamd/tests/src/hipStream.h
+2 -2
View File
@@ -85,7 +85,7 @@ void initArrays(T **Ad, T **Ah,
HIPCHECK( hipMalloc(Ad, NBytes));
}
if(usePinnedHost){
HIPCHECK( hipMallocHost(Ah, NBytes));
HIPCHECK( hipHostAlloc((void**)Ah, NBytes, hipHostAllocDefault));
}
else{
*Ah = new T[N];
@@ -102,7 +102,7 @@ void initArrays(T **Ad, size_t N,
HIPCHECK( hipMalloc(Ad, NBytes));
}else{
if(usePinnedHost){
HIPCHECK(hipMallocHost(Ad, NBytes));
HIPCHECK(hipHostAlloc((void**)Ad, NBytes, hipHostAllocDefault));
}else{
*Ad = new T[N];
HIPASSERT(*Ad != NULL);
hipamd/tests/src/test_common.h
+4 -4
View File
@@ -141,13 +141,13 @@ void initArrays(T **A_d, T **B_d, T **C_d,
if (usePinnedHost) {
if (A_h) {
HIPCHECK ( hipMallocHost(A_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)A_h, Nbytes, hipHostAllocDefault) );
}
if (B_h) {
HIPCHECK ( hipMallocHost(B_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)B_h, Nbytes, hipHostAllocDefault) );
}
if (C_h) {
HIPCHECK ( hipMallocHost(C_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)C_h, Nbytes, hipHostAllocDefault) );
}
} else {
if (A_h) {
@@ -258,7 +258,7 @@ struct Pinned {
static void *Alloc(size_t sizeBytes)
{
void *p;
HIPCHECK(hipMallocHost(&p, sizeBytes));
HIPCHECK(hipHostAlloc((void**)&p, sizeBytes, hipHostAllocDefault));
return p;
};
};