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

Re-sync with upstream.

Browse Source
This commit is contained in:
Alex Voicu
2017-11-20 15:34:50 +00:00
parent 45ac5b70df a1e0c96689
commit 30d90dab38
65 changed files with 320 additions and 265 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
+1 -1
View File
@@ -34,7 +34,7 @@ THE SOFTWARE.
#define NUM_STREAMS 2
__global__ void Iter(hipLaunchParm lp, int *Ad, int num){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
// Kernel loop designed to execute very slowly... ... ... so we can test timing-related behavior below
if(tx == 0){
for(int i = 0; i<num;i++){
hipamd/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
+1 -1
View File
@@ -33,7 +33,7 @@ THE SOFTWARE.
#define SIZE LEN*sizeof(float)
__global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Cd[tx] = Ad[tx] + Bd[tx];
}
hipamd/tests/src/runtimeApi/memory/hipHostMalloc.cpp
+13 -13
View File
@@ -33,13 +33,13 @@
#define SIZE LEN*sizeof(float)
__global__ void Add(float *Ad, float *Bd, float *Cd){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Cd[tx] = Ad[tx] + Bd[tx];
}
__global__ void Set(int *Ad, int val){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tx] = val;
}
@@ -52,13 +52,13 @@ std::vector<std::string> syncMsg = {"event", "stream", "device"};
void CheckHostPointer(int numElements, int *ptr, unsigned eventFlags, int syncMethod, std::string msg)
{
std::cerr << "test: CheckHostPointer " << msg
std::cerr << "test: CheckHostPointer " << msg
//<< " HIP_COHERENT_HOST_ALLOC=" << HIP_COHERENT_HOST_ALLOC
//<< " HIP_EVENT_SYS_RELEASE=" << HIP_EVENT_SYS_RELEASE
<< " eventFlags = " << std::hex << eventFlags
<< ((eventFlags & hipEventReleaseToDevice) ? " hipEventReleaseToDevice" : "")
<< ((eventFlags & hipEventReleaseToSystem) ? " hipEventReleaseToSystem" : "")
<< " ptr=" << ptr
<< " eventFlags = " << std::hex << eventFlags
<< ((eventFlags & hipEventReleaseToDevice) ? " hipEventReleaseToDevice" : "")
<< ((eventFlags & hipEventReleaseToSystem) ? " hipEventReleaseToSystem" : "")
<< " ptr=" << ptr
<< " syncMethod=" << syncMsg[syncMethod] << "\n";
hipStream_t s;
@@ -93,7 +93,7 @@ void CheckHostPointer(int numElements, int *ptr, unsigned eventFlags, int syncMe
default:
assert(0);
};
for (int i=0; i<numElements; i++) {
if (ptr[i] != expected) {
printf ("mismatch at %d: %d != %d\n", i, ptr[i], expected);
@@ -153,7 +153,7 @@ int main(){
size_t sizeBytes = numElements * sizeof (int);
#ifdef __HIP_PLATFORM_HCC__
{
{
// Stimulate error condition:
int *A = &numElements;
HIPCHECK_API(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent|hipHostMallocNonCoherent), hipErrorInvalidValue);
@@ -174,7 +174,7 @@ int main(){
// agent-scope releases don't provide host visibility, don't use them here:
}
if (1) {
if (1) {
int *A = nullptr;
HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
const char *ptrType = "coherent";
@@ -189,14 +189,14 @@ int main(){
// Check defaults:
if (1) {
if (1) {
int *A = nullptr;
HIPCHECK(hipHostMalloc((void**)&A, sizeBytes));
const char *ptrType = "default";
CheckHostPointer(numElements, A, 0, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_EVENT, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_EVENT, ptrType);
@@ -206,7 +206,7 @@ int main(){
}
passed();
}
hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
+6 -6
View File
@@ -29,13 +29,13 @@ THE SOFTWARE.
#include<malloc.h>
__global__ void Inc(hipLaunchParm lp, float *Ad){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tx] = Ad[tx] + float(1);
}
template<typename T>
void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internalRegister)
void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internalRegister)
{
A = A + offset;
numElements -= offset;
@@ -56,7 +56,7 @@ void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internal
HIPCHECK(hipMemset(Bd, 13.0f, sizeBytes));
//
//
HIPCHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));
@@ -81,7 +81,7 @@ int main(int argc, char *argv[])
const size_t size = N * sizeof(float);
if (p_tests & 0x1) {
if (p_tests & 0x1) {
float *A, **Ad;
int num_devices;
HIPCHECK(hipGetDeviceCount(&num_devices));
@@ -115,7 +115,7 @@ int main(int argc, char *argv[])
}
if (p_tests & 0x6) {
if (p_tests & 0x6) {
// Sensitize HIP bug if device does not match where the memory was registered.
HIPCHECK(hipSetDevice(0));
@@ -129,7 +129,7 @@ int main(int argc, char *argv[])
HIPCHECK(hipMalloc(&Bd, size));
// TODO - set to 128
#define OFFSETS_TO_TRY 128
#define OFFSETS_TO_TRY 128
assert (N>OFFSETS_TO_TRY);
if (p_tests & 0x2) {
hipamd/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp
+12 -12
View File
@@ -59,7 +59,7 @@ struct HostTraits<Pinned>
static const char *Name() { return "Pinned"; } ;
static void *Alloc(size_t sizeBytes) {
void *p;
void *p;
HIPCHECK(hipHostMalloc((void**)&p, sizeBytes, hipHostMallocDefault));
return p;
};
@@ -67,11 +67,11 @@ struct HostTraits<Pinned>
template<typename T>
__global__ void
__global__ void
addK (hipLaunchParm lp, T *A, T K, size_t numElements)
{
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x ;
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
for (size_t i=offset; i<numElements; i+=stride) {
A[i] = A[i] + K;
@@ -85,7 +85,7 @@ addK (hipLaunchParm lp, T *A, T K, size_t numElements)
//IN: numInflight : number of copies inflight at any time:
//IN: numPongs = number of iterations to run (iteration)
template<typename T, class AllocType>
void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs, bool doHostSide)
void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs, bool doHostSide)
{
HIPASSERT(numElements % numInflight == 0); // Must be evenly divisible.
size_t Nbytes = numElements*sizeof(T);
@@ -95,7 +95,7 @@ void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
printf ("-----------------------------------------------------------------------------------------------\n");
printf ("testing: %s<%s> Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
printf ("testing: %s<%s> Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
__func__, HostTraits<AllocType>::Name(), Nbytes, (double)(Nbytes)/1024.0/1024.0, numPongs, numInflight, eachCopyElements, eachCopyBytes);
T *A_h = NULL;
@@ -176,7 +176,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
size_t eachCopyBytes = eachCopyElements * sizeof(T);
printf ("-----------------------------------------------------------------------------------------------\n");
printf ("testing: %s Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
printf ("testing: %s Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
__func__, Nbytes, (double)(Nbytes)/1024.0/1024.0, numCopies, eachCopyElements, eachCopyBytes);
T *A_d;
@@ -194,7 +194,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
//stream=0; // fixme TODO
for (int i=0; i<numCopies; i++)
for (int i=0; i<numCopies; i++)
{
HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h1[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
@@ -204,7 +204,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
HIPCHECK(hipDeviceSynchronize());
}
for (int i=0; i<numCopies; i++)
for (int i=0; i<numCopies; i++)
{
HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
HIPCHECK(hipMemcpyAsync(&A_h2[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
@@ -252,7 +252,7 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
hipStream_t *stream = (hipStream_t*)malloc(sizeof(hipStream_t) * nStreams);
if (useNullStream) {
if (useNullStream) {
nStreams = 1;
stream[0] = NULL;
} else {
@@ -262,7 +262,7 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
}
size_t workLeft = N;
size_t workLeft = N;
size_t workPerStream = N / nStreams;
for (int i = 0; i < nStreams; ++i) {
size_t work = (workLeft < workPerStream) ? workLeft : workPerStream;
@@ -287,7 +287,7 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
} else {
HIPCHECK ( hipMemcpyAsync(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost, stream[i]));
}
}
}
HIPCHECK (hipDeviceSynchronize());
hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
+1 -1
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* HIT_END
*/
hipamd/tests/src/runtimeApi/memory/hipMemoryAllocateCoherent.cpp
+1 -1
View File
@@ -31,7 +31,7 @@ THE SOFTWARE.
__global__ void Kernel(hipLaunchParm lp,volatile float* hostRes)
{
int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hostRes[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
+6 -6
View File
@@ -24,7 +24,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
* RUN: %t EXCLUDE_HIP_PLATFORM all
* RUN: %t EXCLUDE_HIP_PLATFORM all
* HIT_END
*/
@@ -63,8 +63,8 @@ int enablePeers(int dev0, int dev1)
__global__ void
memsetIntKernel(int * ptr, const int val, size_t numElements)
{
int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
int stride = hipBlockDim_x * hipGridDim_x ;
int gid = (blockIdx.x * blockDim.x + threadIdx.x);
int stride = blockDim.x * gridDim.x ;
for (size_t i= gid; i< numElements; i+=stride){
ptr[i] = val;
}
@@ -73,15 +73,15 @@ memsetIntKernel(int * ptr, const int val, size_t numElements)
__global__ void
memcpyIntKernel(const int * src, int* dst, size_t numElements)
{
int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
int stride = hipBlockDim_x * hipGridDim_x ;
int gid = (blockIdx.x * blockDim.x + threadIdx.x);
int stride = blockDim.x * gridDim.x ;
for (size_t i= gid; i< numElements; i+=stride){
dst[i] = src[i];
}
};
// CHeck arrays in reverse order, to more easily detect cases where
// CHeck arrays in reverse order, to more easily detect cases where
// the copy is "partially" done.
void checkReverse(const int *ptr, int numElements, int expected) {
for (int i=numElements-1; i>=0; i--) {
hipamd/tests/src/runtimeApi/module/hipModule.cpp
+1 -1
View File
@@ -34,7 +34,7 @@ THE SOFTWARE.
#define kernel_name "hello_world"
__global__ void Cpy(hipLaunchParm lp, float *Ad, float* Bd){
int tx = hipThreadIdx_x;
int tx = threadIdx.x;
Bd[tx] = Ad[tx];
}
hipamd/tests/src/runtimeApi/module/vcpy_kernel.cpp
+1 -1
View File
@@ -24,7 +24,7 @@ THE SOFTWARE.
extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
{
int tx = hipThreadIdx_x;
int tx = threadIdx.x;
b[tx] = a[tx];
}
hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
+3 -3
View File
@@ -35,7 +35,7 @@ THE SOFTWARE.
template<typename T>
__global__ void Inc(hipLaunchParm lp, T *Array){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Array[tx] = Array[tx] + T(1);
}
@@ -116,7 +116,7 @@ int main(int argc, char **argv)
}
const size_t size = N * sizeof(float);
for (int i=0; i< iterations; i++) {
std::thread t1(run1, size, stream[0]);
@@ -126,7 +126,7 @@ int main(int argc, char **argv)
// std::cout<<"T1"<<std::endl;
t2.join();
// std::cout<<"T2"<<std::endl;
t3.join();
t3.join();
}
passed();
}
hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp
+2 -2
View File
@@ -29,7 +29,7 @@ THE SOFTWARE.
const int NN = 1 << 21;
__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
int tid = hipThreadIdx_x;
int tid = threadIdx.x;
if(tid < 1){
for(int i=0;i<n;i++){
x[i] = sqrt(powf(3.14159,i));
@@ -39,7 +39,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
}
__global__ void nKernel(hipLaunchParm lp, float *y){
int tid = hipThreadIdx_x;
int tid = threadIdx.x;
y[tid] = y[tid] + 1.0f;
}
hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp
+2 -2
View File
@@ -31,7 +31,7 @@ THE SOFTWARE.
const int NN = 1 << 21;
__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
int tid = hipThreadIdx_x;
int tid = threadIdx.x;
if(tid < 1){
for(int i=0;i<n;i++){
x[i] = sqrt(powf(3.14159,i));
@@ -41,7 +41,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
}
__global__ void nKernel(hipLaunchParm lp, float *y){
int tid = hipThreadIdx_x;
int tid = threadIdx.x;
y[tid] = y[tid] + 1.0f;
}
hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp
+2 -2
View File
@@ -41,8 +41,8 @@ vectorADDRepeat(hipLaunchParm lp,
size_t NELEM,
int repeat)
{
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x ;
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
for (int j=1; j<=repeat;j++) {
for (size_t i=offset; i<NELEM; i+=stride) {
hipamd/tests/src/runtimeApi/stream/hipStream.h
+4 -4
View File
@@ -73,7 +73,7 @@ void D2H(T *Dst, T *Src, size_t size){
template<typename T>
__global__ void Inc(hipLaunchParm lp, T *In){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
In[tx] = In[tx] + 1;
}
@@ -94,12 +94,12 @@ void initArrays(T **Ad, T **Ah,
}
template<typename T>
void initArrays(T **Ad, size_t N,
bool deviceMemory = false,
void initArrays(T **Ad, size_t N,
bool deviceMemory = false,
bool usePinnedHost = false){
size_t NBytes = N * sizeof(T);
if(deviceMemory){
HIPCHECK( hipMalloc(Ad, NBytes));
HIPCHECK( hipMalloc(Ad, NBytes));
}else{
if(usePinnedHost){
HIPCHECK(hipHostMalloc((void**)Ad, NBytes, hipHostMallocDefault));
hipamd/tests/src/runtimeApi/synchronization/copy_coherency.cpp
+20 -20
View File
@@ -27,14 +27,14 @@ THE SOFTWARE.
// TODO - add code object support here.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
* RUN: %t
* RUN: %t
* HIT_END
*/
// Test cache management (fences) and synchronization between kernel and copy commands.
// Exhaustively tests 3 command types (copy, kernel, module kernel),
// many sync types (see SyncType), followed by another command, across a sweep
// Exhaustively tests 3 command types (copy, kernel, module kernel),
// many sync types (see SyncType), followed by another command, across a sweep
// of data sizes designed to stress various levels of the memory hierarchy.
#include "hip/hip_runtime.h"
@@ -102,8 +102,8 @@ MemcpyFunction g_moduleMemcpy("memcpyInt.hsaco", "memcpyIntKernel");
__global__ void
memsetIntKernel(int * ptr, const int val, size_t numElements)
{
int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
int stride = hipBlockDim_x * hipGridDim_x ;
int gid = (blockIdx.x * blockDim.x + threadIdx.x);
int stride = blockDim.x * gridDim.x ;
for (size_t i= gid; i< numElements; i+=stride){
ptr[i] = val;
}
@@ -112,15 +112,15 @@ memsetIntKernel(int * ptr, const int val, size_t numElements)
__global__ void
memcpyIntKernel(int *dst, const int * src, size_t numElements)
{
int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
int stride = hipBlockDim_x * hipGridDim_x ;
int gid = (blockIdx.x * blockDim.x + threadIdx.x);
int stride = blockDim.x * gridDim.x ;
for (size_t i= gid; i< numElements; i+=stride){
dst[i] = src[i];
}
};
// CHeck arrays in reverse order, to more easily detect cases where
// CHeck arrays in reverse order, to more easily detect cases where
// the copy is "partially" done.
void checkReverse(const int *ptr, int numElements, int expected) {
int mismatchCnt = 0;
@@ -157,7 +157,7 @@ const char * CmdTypeStr(CmdType c)
ENUM_CASE_STR(MODULE_KERNEL);
default: return "UNKNOWN";
};
}
}
enum SyncType {
@@ -223,16 +223,16 @@ void resetInputs( int * Ad, int * Bd, int *Cd, int *Ch, size_t numElements, int
// Intended to test proper synchronization and cache flushing between CMDA and CMDB.
// CMD are of type CmdType. All command copy memory, using either hipMemcpyAsync or kernel implementations.
// CmdA copies from Ad to Bd,
// CmdA copies from Ad to Bd,
// Some form of synchronization is applied.
// Then cmdB copies from Bd to Cd.
//
// Cd is then copied to host Ch using a memory copy.
//
// Correct result at the end is that Ch contains the contents originally in Ad (integer 0x42)
void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType,
void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType,
hipStream_t stream1, hipStream_t stream2, int numElements,
int * Ad, int * Bd, int *Cd, int *Ch,
int * Ad, int * Bd, int *Cd, int *Ch,
int expected)
{
hipEvent_t e;
@@ -241,14 +241,14 @@ void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType,
resetInputs(Ad, Bd, Cd, Ch, numElements, expected);
const size_t sizeElements = numElements * sizeof(int);
fprintf (stderr, "test: runTest with %zu bytes (%6.2f MB) cmdA=%s; sync=%s; cmdB=%s\n",
fprintf (stderr, "test: runTest with %zu bytes (%6.2f MB) cmdA=%s; sync=%s; cmdB=%s\n",
sizeElements, (double) (sizeElements/1024.0), CmdTypeStr(cmdAType), SyncTypeStr(syncType), CmdTypeStr(cmdBType));
if (SKIP_MODULE_KERNEL && ((cmdAType == MODULE_KERNEL) || (cmdBType == MODULE_KERNEL))) {
fprintf (stderr, "warn: skipping since test infra does not yet support modules\n");
return;
}
// Step A:
runCmd(cmdAType, Bd, Ad, stream1, numElements);
@@ -334,7 +334,7 @@ void testWrapper(size_t numElements)
fprintf (stderr, "test: init complete, start running tests\n");
runTestImpl(COPY, EVENT_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, EVENT_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
for (int cmdA=0; cmdA<MAX_CmdType; cmdA++) {
for (int cmdB=0; cmdB<MAX_CmdType; cmdB++) {
@@ -347,7 +347,7 @@ void testWrapper(size_t numElements)
//case STREAM_QUERY:
case STREAM_SYNC:
case DEVICE_SYNC:
runTestImpl(CmdType(cmdA), SyncType(syncMode), CmdType(cmdB), stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(CmdType(cmdA), SyncType(syncMode), CmdType(cmdB), stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
break;
default:
break;
@@ -358,11 +358,11 @@ void testWrapper(size_t numElements)
}
#if 0
runTestImpl(COPY, STREAM_SYNC, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_WAIT_EVENT, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_SYNC, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_WAIT_EVENT, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_WAIT_EVENT, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
runTestImpl(COPY, STREAM_WAIT_EVENT, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
#endif
HIPCHECK(hipFree(Ad));
hipamd/tests/src/runtimeApi/synchronization/memcpyInt.device.cpp
+3 -3
View File
@@ -1,12 +1,12 @@
#include <hip/hip_runtime.h>
extern "C" __global__ void
memcpyIntKernel(hipLaunchParm lp, int *dst, const int * src, size_t numElements)
{
int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
int stride = hipBlockDim_x * hipGridDim_x ;
int gid = (blockIdx.x * blockDim.x + threadIdx.x);
int stride = blockDim.x * gridDim.x ;
for (size_t i= gid; i< numElements; i+=stride){
dst[i] = src[i];
}