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

Merge branch 'amd-master' into amd-master-next

Browse Source 
Change-Id: I3094c15008093f2072bcd38aca4ea90aeae2d97b
This commit is contained in:
Maneesh Gupta
2020-04-07 06:57:42 -04:00
parent a48b312aa9
commit 2af31479e2
67 changed files with 2278 additions and 601 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/Negative/memory/hipMemcpyFromSymbol.cpp
+46
View File
@@ -0,0 +1,46 @@
/*
Copyright (c) 2015-Present 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.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE 1024
int main(){
void *Sd;
hipError_t e;
char S[SIZE]="This is not a device symbol";
HIPCHECK(hipMalloc(&Sd,SIZE));
e = hipMemcpyFromSymbol(S, HIP_SYMBOL(Sd), SIZE, 0, hipMemcpyDeviceToHost);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyFromSymbol(S, NULL, SIZE, 0, hipMemcpyDeviceToHost);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
passed();
}
tests/src/Negative/memory/hipMemcpyFromSymbolAsync.cpp
+49
View File
@@ -0,0 +1,49 @@
/*
Copyright (c) 2015-Present 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.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE 1024
int main(){
void *Sd;
hipError_t e;
char S[SIZE]="This is not a device symbol";
HIPCHECK(hipMalloc(&Sd,SIZE));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
e = hipMemcpyFromSymbolAsync(S, HIP_SYMBOL(Sd), SIZE, 0, hipMemcpyDeviceToHost, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyFromSymbolAsync(S, NULL, SIZE, 0, hipMemcpyDeviceToHost, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
passed();
}
tests/src/Negative/memory/hipMemcpyToSymbol.cpp
+46
View File
@@ -0,0 +1,46 @@
/*
Copyright (c) 2015-Present 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.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE 1024
int main(){
void *Sd;
hipError_t e;
char S[SIZE]="This is not a device symbol";
HIPCHECK(hipMalloc(&Sd,SIZE));
e = hipMemcpyToSymbol(HIP_SYMBOL(Sd), S, SIZE, 0, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyToSymbol(NULL, S, SIZE, 0, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
passed();
}
tests/src/Negative/memory/hipMemcpyToSymbolAsync.cpp
+49
View File
@@ -0,0 +1,49 @@
/*
Copyright (c) 2015-Present 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.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE 100
int main(){
void *Sd;
hipError_t e;
char S[SIZE]="This is not a device symbol";
HIPCHECK(hipMalloc(&Sd,SIZE));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
e = hipMemcpyToSymbolAsync(HIP_SYMBOL(Sd), S, SIZE, 0, hipMemcpyHostToDevice, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
e = hipMemcpyToSymbolAsync(NULL, S, SIZE, 0, hipMemcpyHostToDevice, stream);
HIPASSERT(e==hipErrorInvalidSymbol);
HIPCHECK(hipFree(Sd));
passed();
}
tests/src/Negative/memory/hipMemory.cpp
+43
View File
@@ -0,0 +1,43 @@
/*
Copyright (c) 2015-Present 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.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM vdi
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE 100
int main(){
hipError_t e;
char str[SIZE]="Hi, I am Ellesemere. What is ur name?";
e = hipMemcpy(0, str, SIZE, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidValue);
e = hipMemcpy(NULL, str, SIZE, hipMemcpyHostToDevice);
HIPASSERT(e==hipErrorInvalidValue);
e = hipMemset(0,99,80);
HIPASSERT(e==hipErrorInvalidValue);
passed();
}
tests/src/Negative/stream/hipStreamCreateWithFlags.cpp
+40
View File
@@ -0,0 +1,40 @@
/*
Copyright (c) 2015-Present 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.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM vdi
* TEST: %t
* HIT_END
*/
#include "test_common.h"
int main(){
hipError_t e;
hipStream_t stream;
e = hipStreamCreateWithFlags(&stream, -1);
HIPASSERT(e==hipErrorInvalidValue);
e = hipStreamCreateWithFlags(&stream, 2);
HIPASSERT(e==hipErrorInvalidValue);
passed();
}
tests/src/deviceLib/hipTestHalf.cpp
+39
View File
@@ -96,6 +96,18 @@ void kernel_hisinf(__half* input, int* output) {
output[tx] = __hisinf(input[tx]);
}
__global__ void testHalfAbs(float* p) {
auto a = __float2half(*p);
a = __habs(a);
*p = __half2float(a);
}
__global__ void testHalf2Abs(float2* p) {
auto a = __float22half2_rn(*p);
a = __habs2(a);
*p = __half22float2(a);
}
#endif
@@ -237,6 +249,31 @@ void checkFunctional() {
return;
}
void checkHalfAbs() {
{
float *p;
hipMalloc(&p, sizeof(float));
float pp = -2.1f;
hipMemcpy(p, &pp, sizeof(float), hipMemcpyDefault);
hipLaunchKernelGGL(testHalfAbs, 1, 1, 0, 0, p);
hipMemcpy(&pp, p, sizeof(float), hipMemcpyDefault);
hipFree(p);
if(pp < 0.0f) { failed("Half Abs failed"); }
}
{
float2 *p;
hipMalloc(&p, sizeof(float2));
float2 pp;
pp.x = -2.1f;
pp.y = -1.1f;
hipMemcpy(p, &pp, sizeof(float2), hipMemcpyDefault);
hipLaunchKernelGGL(testHalf2Abs, 1, 1, 0, 0, p);
hipMemcpy(&pp, p, sizeof(float2), hipMemcpyDefault);
hipFree(p);
if(pp.x < 0.0f || pp.y < 0.0f) { failed("Half2 Abs Test Failed"); }
}
}
int main() {
bool* result{nullptr};
hipMemAllocHost((void**)&result, sizeof(result));
@@ -260,5 +297,7 @@ int main() {
// run some functional checks
checkFunctional();
checkHalfAbs();
passed();
}
tests/src/deviceLib/hip_floatnTM.cpp
+239
View File
@@ -0,0 +1,239 @@
/*
Copyright (c) 2015-2019 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.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc HIPCC_OPTIONS -std=c++14
* TEST: %t
* HIT_END
*/
#include <hip/hip_runtime.h>
#include <type_traits>
#include <random>
#include "test_common.h"
static std::random_device dev;
static std::mt19937 rng(dev());
template <typename T, typename M>
__host__ __device__ inline constexpr int count() {
return sizeof(T) / sizeof(M);
}
inline float getRandomFloat(float min = 10, float max = 100) {
std::uniform_real_distribution<float> gen(min, max);
return gen(rng);
}
template <typename T, typename B>
void fillMatrix(T* a, int size) {
for (int i = 0; i < size; i++) {
T t;
t.x = getRandomFloat();
if constexpr (count<T, B>() >= 2) t.y = getRandomFloat();
if constexpr (count<T, B>() >= 3) t.z = getRandomFloat();
if constexpr (count<T, B>() >= 4) t.w = getRandomFloat();
a[i] = t;
}
}
// Test operations
template <typename T, typename B>
__host__ __device__ void testOperations(T& a, T& b) {
a.x += b.x;
a.x++;
b.x++;
if constexpr (count<T, B>() >= 2) {
a.y = b.x;
a.x = b.y;
}
if constexpr (count<T, B>() >= 3) {
if (a.x > 0) b.x /= a.x;
a.x *= b.z;
a.y--;
}
if constexpr (count<T, B>() >= 4) {
b.w = a.x;
a.w += (-b.y);
}
}
template <typename T, typename B>
__global__ void testOperationsGPU(T* d_a, T* d_b, int size) {
int id = threadIdx.x;
if (id > size) return;
T &a = d_a[id];
T &b = d_b[id];
testOperations<T, B>(a, b);
}
template <typename T>
void dcopy(T* a, T* b, int size) {
for (int i = 0; i < size; i++) {
a[i] = b[i];
}
}
template <typename T>
bool isEqual(T* a, T* b, int size) {
for (int i = 0; i < size; i++) {
if (a[i] != b[i]) {
return false;
}
}
return true;
}
// Main function that tests type
// T = what you want to test
// D = pack of 1 i.e. float1 int1
template <typename T, typename D>
void testType(int msize) {
T *fa, *fb, *fc, *h_fa, *h_fb;
fa = new T[msize];
fb = new T[msize];
fc = new T[msize];
h_fa = new T[msize];
h_fb = new T[msize];
T *d_fa, *d_fb;
constexpr int c = count<T, D>();
if (c <= 0 || c >= 5) {
failed("Invalid Size\n");
}
fillMatrix<T, D>(fa, msize);
dcopy(fb, fa, msize);
dcopy(h_fa, fa, msize);
dcopy(h_fb, fa, msize);
for (int i = 0; i < msize; i++) testOperations<T, D>(h_fa[i], h_fb[i]);
hipMalloc(&d_fa, sizeof(T) * msize);
hipMalloc(&d_fb, sizeof(T) * msize);
hipMemcpy(d_fa, fa, sizeof(T) * msize, hipMemcpyHostToDevice);
hipMemcpy(d_fb, fb, sizeof(T) * msize, hipMemcpyHostToDevice);
auto kernel = testOperationsGPU<T, D>;
hipLaunchKernelGGL(kernel, 1, msize, 0, 0, d_fa, d_fb, msize);
hipMemcpy(fc, d_fa, sizeof(T) * msize, hipMemcpyDeviceToHost);
bool pass = true;
if (!isEqual<T>(h_fa, fc, msize)) {
pass = false;
}
delete[] fa;
delete[] fb;
delete[] fc;
delete[] h_fa;
delete[] h_fb;
hipFree(d_fa);
hipFree(d_fb);
if (!pass) {
failed("Failed");
}
}
int main() {
const int msize = 100;
// double
testType<double1, double1>(msize);
testType<double2, double1>(msize);
testType<double3, double1>(msize);
testType<double4, double1>(msize);
// floats
testType<float1, float1>(msize);
testType<float2, float1>(msize);
testType<float3, float1>(msize);
testType<float4, float1>(msize);
// ints
testType<int1, int1>(msize);
testType<int2, int1>(msize);
testType<int3, int1>(msize);
testType<int4, int1>(msize);
// chars
testType<char1, char1>(msize);
testType<char2, char1>(msize);
testType<char3, char1>(msize);
testType<char4, char1>(msize);
// long
testType<long1, long1>(msize);
testType<long2, long1>(msize);
testType<long3, long1>(msize);
testType<long4, long1>(msize);
// longlong
testType<longlong1, longlong1>(msize);
testType<longlong2, longlong1>(msize);
testType<longlong3, longlong1>(msize);
testType<longlong4, longlong1>(msize);
// short
testType<short1, short1>(msize);
testType<short2, short1>(msize);
testType<short3, short1>(msize);
testType<short4, short1>(msize);
// uints
testType<uint1, uint1>(msize);
testType<uint2, uint1>(msize);
testType<uint3, uint1>(msize);
testType<uint4, uint1>(msize);
// uchars
testType<uchar1, uchar1>(msize);
testType<uchar2, uchar1>(msize);
testType<uchar3, uchar1>(msize);
testType<uchar4, uchar1>(msize);
// ulong
testType<ulong1, ulong1>(msize);
testType<ulong2, ulong1>(msize);
testType<ulong3, ulong1>(msize);
testType<ulong4, ulong1>(msize);
// ulonglong
testType<ulonglong1, ulonglong1>(msize);
testType<ulonglong2, ulonglong1>(msize);
testType<ulonglong3, ulonglong1>(msize);
testType<ulonglong4, ulonglong1>(msize);
// ushort
testType<ushort1, ushort1>(msize);
testType<ushort2, ushort1>(msize);
testType<ushort3, ushort1>(msize);
testType<ushort4, ushort1>(msize);
passed();
}
tests/src/hiprtc/hiprtcGetLoweredName.cpp
+1 -1
View File
@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc vdi
* TEST: %t
* HIT_END
*/
tests/src/hiprtc/saxpy.cpp
+1 -1
View File
@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
* BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc vdi
* TEST: %t
* HIT_END
*/
tests/src/runtimeApi/memory/hipMemcpyNegetiveTests.cpp
+53
View File
@@ -0,0 +1,53 @@
/*
* Copyright (c) 2019-2020 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 WARRANNTY OF ANY KIND, EXPRESS OR
* IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
* THE SOFTWARE.
* */
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc vdi
* TEST: %t
* HIT_END
*/
#include "test_common.h"
int main() {
int* A;
int* Ad;
int* Bd;
// Allocation
HIPCHECK(hipMalloc((void**)&Ad, sizeof(int)));
HIPCHECK(hipMalloc((void**)&Bd, sizeof(int)));
HIPCHECK(hipHostMalloc((void**)&A,sizeof(int)));
// Kind should be ignored and test should pass even for incorrect kind
HIPCHECK(hipMemcpy(Ad, A, sizeof(int), hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(A, Ad, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Ad, Bd, sizeof(int), hipMemcpyHostToHost));
HIPCHECK(hipMemcpy(A, A, sizeof(int), hipMemcpyDeviceToDevice));
// nullptr passed as source or destination pointer
HIPASSERT(hipSuccess != hipMemcpy(nullptr, A, sizeof(int), hipMemcpyHostToDevice));
HIPASSERT(hipSuccess != hipMemcpy(Ad, nullptr, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipFree(Ad));
HIPCHECK(hipFree(Bd));
HIPCHECK(hipFree(A));
passed();
}
tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp
+2 -2
View File
@@ -107,8 +107,8 @@ void run(const std::vector<char>& buffer) {
hipFree(Ad);
hipFree(Bd);
delete A;
delete B;
delete[] A;
delete[] B;
hipCtxDestroy(context);
}
tests/src/runtimeApi/stream/StreamAddCallback.cpp
+145
View File
@@ -0,0 +1,145 @@
#include <stdio.h>
#include <hip/hip_runtime.h>
#include <unistd.h>
#include "test_common.h"
#include <atomic>
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
enum class ExecState
{
EXEC_NOT_STARTED,
EXEC_STARTED,
EXEC_CB_STARTED,
EXEC_CB_FINISHED,
EXEC_FINISHED
};
struct UserData
{
size_t size;
int* ptr;
};
// Global variable to check exection order
std::atomic<ExecState> gData(ExecState::EXEC_NOT_STARTED);
void myCallback(hipStream_t stream, hipError_t status, void* user_data)
{
if(gData.load() != ExecState::EXEC_STARTED)
return; // Error hence return early
gData.store(ExecState::EXEC_CB_STARTED);
UserData* data = reinterpret_cast<UserData*>(user_data);
printf("Callback started\n");
sleep(1);
printf("Callback ending.\n");
gData.store(ExecState::EXEC_CB_FINISHED);
}
bool test(int count)
{
printf("\n============ Test iteration %d =============\n",count);
// Stream
hipStream_t stream;
bool result = true;
gData.store(ExecState::EXEC_STARTED);
HIPCHECK(hipStreamCreate(&stream));
// Array size
size_t size = 10000;
// Device array
int *data = NULL;
HIPCHECK(hipMalloc((void**)&data, sizeof(int) * size));
// Initialize device array to -1
HIPCHECK(hipMemset(data, -1, sizeof(int) * size));
// Host array
int *host = NULL;
HIPCHECK(hipHostMalloc((void**)&host, sizeof(int) * size));
// Print host ptr address
printf("In main thread\n");
// Initialize user_data for callback
UserData arg;
arg.size = size;
arg.ptr = host;
// Synchronize device
HIPCHECK(hipDeviceSynchronize());
// Asynchronous copy from device to host
HIPCHECK(hipMemcpyAsync(host, data, sizeof(int) * size, hipMemcpyDeviceToHost, stream));
// Asynchronous memset on device
HIPCHECK(hipMemsetAsync(data, 0, sizeof(int) * size, stream));
// Add callback - should happen after hipMemsetAsync()
HIPCHECK(hipStreamAddCallback(stream, myCallback, &arg, 0));
printf("Will wait in main thread until callback completes\n");
//This should synchronize the stream (including the callback)
HIPCHECK(hipStreamSynchronize(stream));
if(gData.load() != ExecState::EXEC_CB_FINISHED)
{
std::cout<<"Callback is not finished\n";
return false;
}
printf("Callback completed will resume main thread execution\n");
if(host[size/2] != -1)
{
// Print some host data that just got copied
printf("Pseudo host data printing (should be -1): %d\n", host[size/2]);
result = false;
}
HIPCHECK(hipMemcpy(host, data, sizeof(int)*size, hipMemcpyDeviceToHost));
if(host[size-1] != 0)
{
printf("Pseudo host data printing (should be 0): %d\n", host[size-1]);
result = false;
}
HIPCHECK(hipFree(data));
HIPCHECK(hipHostFree(host));
HIPCHECK(hipStreamDestroy(stream));
gData.store(ExecState::EXEC_FINISHED);
return result;
}
int main()
{
// Test involves multithreading hence running multiple times
// to make sure consitency in the behavior
bool status = true;
for(int i=0; i < 10; i++){
status = test(i+1);
if(status == false)
{
failed("Test Failed!\n");
break;
}
}
if(status == true) passed();
return 0;
}
tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp
+409
View File
@@ -0,0 +1,409 @@
#include <hip/hip_runtime.h>
#include <stdexcept>
#include <memory>
#include <functional>
#include <mutex>
#include <condition_variable>
#include <thread>
#include <future>
#include "test_common.h"
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM vdi
* TEST: %t
* HIT_END
*/
#define WORKAROUND 0 // Enable (1) this to make stream thread-safe by a workaround
template<bool IsBlocking> // <true> = queue blocks, until task is finished in enqueue(queue,task)
class QueueHipRt;
// Queue types used in the tests
using TestQueues = std::tuple<QueueHipRt<true>, QueueHipRt<false>>;
// --- Implementation
#define HIP_ASSERT(x) (assert((x)==hipSuccess))
#define HIP_ASSERT_IGNORE(x,ign) auto err=x; HIP_ASSERT(err==ign ? hipSuccess : err)
#ifdef __HIP_PLATFORM_HCC__
#define HIPRT_CB
#endif
template<bool isBlocking>
static auto currentThreadWaitFor(QueueHipRt<isBlocking> const & queue) -> void;
template<bool IsBlocking>
class QueueHipRt
{
public:
static constexpr bool isBlocking = IsBlocking;
//-----------------------------------------------------------------------------
QueueHipRt(
int dev) :
m_dev(dev),
m_HipQueue()
{
HIP_ASSERT(
hipSetDevice(
m_dev));
HIP_ASSERT(
hipStreamCreateWithFlags(
&m_HipQueue,
hipStreamNonBlocking));
}
//-----------------------------------------------------------------------------
QueueHipRt(QueueHipRt const &) = delete;
//-----------------------------------------------------------------------------
QueueHipRt(QueueHipRt &&) = delete;
//-----------------------------------------------------------------------------
auto operator=(QueueHipRt const &) -> QueueHipRt & = delete;
//-----------------------------------------------------------------------------
auto operator=(QueueHipRt &&) -> QueueHipRt & = delete;
//-----------------------------------------------------------------------------
~QueueHipRt()
{
if(isBlocking) {
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
// we are a non-blocking queue, so we have to wait here with its destruction until all spawned tasks have been processed
currentThreadWaitFor(*this);
#endif
}
HIP_ASSERT(
hipSetDevice(
m_dev));
HIP_ASSERT(
hipStreamDestroy(
m_HipQueue));
}
public:
int m_dev; //!< The device this queue is bound to.
hipStream_t m_HipQueue;
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
int m_callees = 0;
std::mutex m_mutex;
#endif
};
template<typename TTask>
struct Enqueue
{
//#############################################################################
enum class CallbackState
{
enqueued,
notified,
finished,
};
//#############################################################################
struct CallbackSynchronizationData : public std::enable_shared_from_this<CallbackSynchronizationData>
{
std::mutex m_mutex;
std::condition_variable m_event;
CallbackState state = CallbackState::enqueued;
};
//-----------------------------------------------------------------------------
static void HIPRT_CB hipRtCallback(hipStream_t /*queue*/, hipError_t /*status*/, void *arg)
{
// explicitly copy the shared_ptr so that this method holds the state even when the executing thread has already finished.
const auto pCallbackSynchronizationData = reinterpret_cast<CallbackSynchronizationData*>(arg)->shared_from_this();
// Notify the executing thread.
{
std::unique_lock<std::mutex> lock(pCallbackSynchronizationData->m_mutex);
pCallbackSynchronizationData->state = CallbackState::notified;
}
pCallbackSynchronizationData->m_event.notify_one();
// Wait for the executing thread to finish the task if it has not already finished.
std::unique_lock<std::mutex> lock(pCallbackSynchronizationData->m_mutex);
if(pCallbackSynchronizationData->state != CallbackState::finished)
{
pCallbackSynchronizationData->m_event.wait(
lock,
[pCallbackSynchronizationData](){
return pCallbackSynchronizationData->state == CallbackState::finished;
}
);
}
}
//-----------------------------------------------------------------------------
template<bool isBlocking>
static auto enqueue(
QueueHipRt<isBlocking> & queue,
TTask const & task)
-> void
{
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
{
// thread-safe callee incrementing
std::lock_guard<std::mutex> guard(queue.m_mutex);
queue.m_callees += 1;
}
#endif
auto pCallbackSynchronizationData = std::make_shared<CallbackSynchronizationData>();
// test example: https://github.com/ROCm-Developer-Tools/HIP/blob/roc-1.9.x/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp
HIP_ASSERT(hipStreamAddCallback(
queue.m_HipQueue,
hipRtCallback,
pCallbackSynchronizationData.get(),
0u));
// We start a new std::thread which stores the task to be executed.
// This circumvents the limitation that it is not possible to call HIP methods within the HIP callback thread.
// The HIP thread signals the std::thread when it is ready to execute the task.
// The HIP thread is waiting for the std::thread to signal that it is finished executing the task
// before it executes the next task in the queue (HIP stream).
std::thread t(
[pCallbackSynchronizationData,
task
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
,&queue // requires queue's destructor to wait for all tasks
#endif
](){
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
// thread-safe task execution and callee decrementing
std::lock_guard<std::mutex> guard(queue.m_mutex);
#endif
// If the callback has not yet been called, we wait for it.
{
std::unique_lock<std::mutex> lock(pCallbackSynchronizationData->m_mutex);
if(pCallbackSynchronizationData->state != CallbackState::notified)
{
pCallbackSynchronizationData->m_event.wait(
lock,
[pCallbackSynchronizationData](){
return pCallbackSynchronizationData->state == CallbackState::notified;
}
);
}
task();
// Notify the waiting HIP thread.
pCallbackSynchronizationData->state = CallbackState::finished;
}
pCallbackSynchronizationData->m_event.notify_one();
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
queue.m_callees -= 1;
#endif
}
);
if(isBlocking)
t.join(); // => waiting for task completion
else
t.detach(); // => do not wait for task completion
}
};
//#############################################################################
//! The HIP RT non-blocking queue test trait specialization.
struct Empty
{
//-----------------------------------------------------------------------------
template<bool isBlocking>
static auto empty(
QueueHipRt<isBlocking> const & queue)
-> bool
{
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
return (queue.m_callees==0);
#else
// Query is allowed even for queues on non current device.
hipError_t ret = hipSuccess;
HIP_ASSERT_IGNORE(
ret = hipStreamQuery(
queue.m_HipQueue),
hipErrorNotReady);
return (ret == hipSuccess);
#endif
}
};
template<bool isBlocking>
auto currentThreadWaitFor(QueueHipRt<isBlocking> const & queue) -> void
{
#if WORKAROUND // NOTE: workaround for unwanted nonblocking hip streams for HCC (NVCC streams are blocking)
while(queue.m_callees>0) {
std::this_thread::sleep_for(std::chrono::milliseconds(10u));
}
#else
// Sync is allowed even for queues on non current device.
HIP_ASSERT( hipStreamSynchronize(
queue.m_HipQueue));
#endif
}
// --- Tests
#define TEMPLATE_LIST_TEST_CASE(TestName) \
template<typename TestType> static void TestName (std::atomic<int> &check); \
static int TestName##Runner () { \
std::atomic<int> check{0}; \
TestName< QueueHipRt<true> >(check); \
fprintf(stderr, "After " #TestName " < QueueHipRt<true> > errors=%d\n", check.load()); \
TestName< QueueHipRt<false> >(check); \
fprintf(stderr, "After " #TestName " < QueueHipRt<false> > errors=%d\n", check.load()); \
return check.load(); \
} \
template<typename TestType> static void TestName (std::atomic<int> &check)
// add 1 if a check fails
#define CHECK(result) do{int arg=(!(result)); fprintf(stderr, "Checking " #result " %d\n", arg); check.fetch_add(arg);}while(false)
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( queueIsInitiallyEmpty )
{
TestType queue{0};
CHECK(Empty::empty(queue));
}
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( queueCallbackIsWorking )
{
std::promise<bool> promise;
auto task = [&](){ promise.set_value(true); };
TestType queue{0};
Enqueue<decltype(task)> enqueue;
enqueue.enqueue(
queue,
task
);
CHECK(promise.get_future().get());
}
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( queueWaitShouldWork )
{
bool CallbackFinished = false;
auto task =
[&CallbackFinished]() noexcept
{
std::this_thread::sleep_for(std::chrono::milliseconds(100u));
CallbackFinished = true;
};
TestType queue{0};
Enqueue<decltype(task)> enqueue;
enqueue.enqueue(
queue,
task
);
currentThreadWaitFor(queue);
CHECK(CallbackFinished);
}
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( queueShouldNotBeEmptyWhenLastTaskIsStillExecutingAndIsEmptyAfterProcessingFinished )
{
bool CallbackFinished = false;
TestType queue{0};
auto task = [&queue, &CallbackFinished, &check]() noexcept
{
CHECK(!Empty::empty(queue));
std::this_thread::sleep_for(std::chrono::milliseconds(100u));
CallbackFinished = true;
};
Enqueue<decltype(task)> enqueue;
enqueue.enqueue(
queue,
task
);
// A non-blocking queue will always stay empty because the task has been executed immediately.
if(!TestType::isBlocking)
{
currentThreadWaitFor(queue);
}
CHECK(Empty::empty(queue));
CHECK(CallbackFinished);
}
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( queueShouldNotExecuteTasksInParallel )
{
std::atomic<bool> taskIsExecuting(false);
std::promise<void> firstTaskFinished;
std::future<void> firstTaskFinishedFuture = firstTaskFinished.get_future();
std::promise<void> secondTaskFinished;
std::future<void> secondTaskFinishedFuture = secondTaskFinished.get_future();
TestType queue{0};
std::thread thread1(
[&queue, &taskIsExecuting, &firstTaskFinished, &check]()
{
auto task1 = [&taskIsExecuting, &firstTaskFinished, &check]() noexcept
{
CHECK(!taskIsExecuting.exchange(true));
std::this_thread::sleep_for(std::chrono::milliseconds(100u));
CHECK(taskIsExecuting.exchange(false));
firstTaskFinished.set_value();
};
Enqueue<decltype(task1)> enqueue;
enqueue.enqueue(
queue,
task1
);
});
std::thread thread2(
[&queue, &taskIsExecuting, &secondTaskFinished, &check]()
{
auto task2 = [&taskIsExecuting, &secondTaskFinished, &check]() noexcept
{
CHECK(!taskIsExecuting.exchange(true));
std::this_thread::sleep_for(std::chrono::milliseconds(100u));
CHECK(taskIsExecuting.exchange(false));
secondTaskFinished.set_value();
};
Enqueue<decltype(task2)> enqueue;
enqueue.enqueue(
queue,
task2
);
});
// Both tasks have to be enqueued
thread1.join();
thread2.join();
currentThreadWaitFor(queue);
firstTaskFinishedFuture.get();
secondTaskFinishedFuture.get();
}
#define TESTER(name) do { \
int result = name (); \
fprintf(stderr, #name " %s\n", result?"Errors":"No Errors"); \
if (result) { failed(#name " failed\n"); } \
} while (false)
int main()
{
TESTER(queueIsInitiallyEmptyRunner);
TESTER(queueCallbackIsWorkingRunner);
TESTER(queueWaitShouldWorkRunner);
TESTER(queueShouldNotBeEmptyWhenLastTaskIsStillExecutingAndIsEmptyAfterProcessingFinishedRunner);
TESTER(queueShouldNotExecuteTasksInParallelRunner);
passed();
}
tests/src/texture/hipTex1DFetchCheckModes.cpp
+122
View File
@@ -0,0 +1,122 @@
/*
Copyright (c) 2019 - present 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.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM vdi
* TEST: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
#include "../test_common.h"
#define N 16
#define offset 3
__global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < N)
val[k] = tex1Dfetch<float>(obj, k+offset);
}
int runTest(hipTextureAddressMode, hipTextureFilterMode);
int main(int argc, char **argv) {
int testResult = runTest(hipAddressModeClamp,hipFilterModePoint);
testResult = runTest(hipAddressModeClamp,hipFilterModeLinear);
testResult = runTest(hipAddressModeWrap,hipFilterModePoint);
testResult = runTest(hipAddressModeWrap,hipFilterModeLinear);
if(testResult) {
passed();
} else {
exit(EXIT_FAILURE);
}
}
int runTest(hipTextureAddressMode addressMode, hipTextureFilterMode filterMode) {
int testResult = 1;
hipCtx_t HipContext;
hipDevice_t HipDevice;
int deviceID = 0;
hipDeviceGet(&HipDevice, deviceID);
hipCtxCreate(&HipContext, 0, HipDevice);
// Allocating the required buffer on gpu device
float *texBuf, *texBufOut;
float val[N], output[N];
for (int i = 0; i < N; i++) {
val[i] = i+1;
output[i] = 0.0;
}
HIPCHECK(hipMalloc(&texBuf, N * sizeof(float)));
HIPCHECK(hipMalloc(&texBufOut, N * sizeof(float)));
HIPCHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice));
HIPCHECK(hipMemset(texBufOut, 0, N * sizeof(float)));
hipResourceDesc resDescLinear;
memset(&resDescLinear, 0, sizeof(resDescLinear));
resDescLinear.resType = hipResourceTypeLinear;
resDescLinear.res.linear.devPtr = texBuf;
resDescLinear.res.linear.desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
resDescLinear.res.linear.sizeInBytes = N * sizeof(float);
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = hipReadModeElementType;
texDesc.addressMode[0] = addressMode;
texDesc.addressMode[1] = addressMode;
texDesc.filterMode = filterMode;
texDesc.normalizedCoords = false;
// Creating texture object
hipTextureObject_t texObj = 0;
HIPCHECK(hipCreateTextureObject(&texObj, &resDescLinear, &texDesc, NULL));
dim3 dimBlock(1, 1, 1);
dim3 dimGrid(N , 1, 1);
hipLaunchKernelGGL(tex1dKernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
texBufOut, texObj);
HIPCHECK(hipDeviceSynchronize());
HIPCHECK(hipMemcpy(output, texBufOut, N * sizeof(float), hipMemcpyDeviceToHost));
for (int i = offset; i < N; i++) {
if (output[i-offset] != val[i]) {
testResult = 0;
break;
}
}
if(testResult){
for(int i = N-offset; i < N; i++){
if (output[i] != 0){
testResult = 0;
break;
}
}
}
HIPCHECK(hipDestroyTextureObject(texObj));
HIPCHECK(hipFree(texBuf));
HIPCHECK(hipFree(texBufOut));
return testResult;
}