2
0

SWDEV-388834 - [catch2][dtest] Kernel tests migrated from direct to catch2 (#345)

Change-Id: I8d1d7c6d5db018301cd76f2e38b5997ae91c15db

[ROCm/hip-tests commit: a461ae2fc9]
Este cometimento está contido em:
ROCm CI Service Account
2023-07-08 20:54:42 +05:30
cometido por GitHub
ascendente e93c1dd330
cometimento a1346c114a
10 ficheiros modificados com 983 adições e 1 eliminações
+8 -1
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021 - 2023 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
@@ -100,3 +100,10 @@ THE SOFTWARE.
* This section describes the stream management types & functions of HIP runtime API.
* @}
*/
/**
* @defgroup KernelTest Kernel Functions Management
* @{
* This section describes the various kernel functions invocation.
* @}
*/
+18
Ver ficheiro
@@ -23,7 +23,25 @@ set(TEST_SRC
hipMemFaultStackAllocation.cc
hipLaunchBounds.cc
hipShflTests.cc
hipTestConstant.cc
hipTestGlobalVariable.cc
hipTestMemKernel.cc
launch_bounds.cc
inline_asm_vadd.cc
)
if(UNIX)
set(TEST_SRC ${TEST_SRC}
hipPrintfKernel.cc)
endif()
string(FIND "${OFFLOAD_ARCH_STR}" "gfx1030" RESULT)
if ("${RESULT}" EQUAL "-1")
set(AMD_TEST_SRC
inline_asm_vmac.cc)
endif()
if(HIP_PLATFORM MATCHES "amd")
set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC})
endif()
# only for AMD
if(HIP_PLATFORM MATCHES "amd")
+67
Ver ficheiro
@@ -0,0 +1,67 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
#include <cstring>
#include "../kernel/printf_common.h"
#define HIP_ENABLE_PRINTF
__global__ void run_printf() {
printf("Hello World\n");
}
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check printf function via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/hipPrintfKernel.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_kernel_ChkPrintf") {
int device_count = 0;
CaptureStream capture(stdout);
HIP_CHECK(hipGetDeviceCount(&device_count));
std::string st = "Hello World";
const char * check = st.c_str();
for (int i = 0; i < device_count; ++i) {
HIP_CHECK(hipSetDevice(i));
hipLaunchKernelGGL(run_printf, dim3(1), dim3(1), 0, 0);
HIP_CHECK(hipDeviceSynchronize());
char* data = new char[st.size()];;
std::ifstream CapturedData = capture.getCapturedData();
CapturedData.getline(data, st.size()+1);
int result = strcmp(data, check);
REQUIRE(result == 0);
delete [] data;
}
}
+76
Ver ficheiro
@@ -0,0 +1,76 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
#define LEN 512
#define SIZE 2048
__constant__ int Value[LEN];
static __global__ void Get(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = Value[tid];
}
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check const variable via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/hipTestConstant.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_kernel_chkConstantViaKernel") {
int *A, *B, *Ad;
A = new int[LEN];
B = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = -1 * i;
B[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(Value), A, SIZE, 0,
hipMemcpyHostToDevice));
hipLaunchKernelGGL(Get, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
REQUIRE(A[i] == B[i]);
}
delete [] A;
delete [] B;
HIP_CHECK(hipFree(Ad));
}
@@ -0,0 +1,101 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
#define LEN 512
#define SIZE 2048
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check constant global variable and global array via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/hipTestGlobalVariable.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
__constant__ int ConstantGlobalVar = 123;
static __global__ void kernel(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = ConstantGlobalVar;
}
void runTestConstantGlobalVar() {
int *A, *Ad;
A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
REQUIRE(123 == A[i]);
}
delete [] A;
HIP_CHECK(hipFree(Ad));
}
__device__ int GlobalArray[LEN];
static __global__ void kernelWrite() {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
GlobalArray[tid] = tid;
}
static __global__ void kernelRead(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = GlobalArray[tid];
}
void runTestGlobalArray() {
int *A, *Ad;
A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
hipLaunchKernelGGL(kernelWrite, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0);
hipLaunchKernelGGL(kernelRead, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
REQUIRE(i == A[i]);
}
delete [] A;
HIP_CHECK(hipFree(Ad));
}
TEST_CASE("Unit_kernel_chkGlobalArrAndGlobalVaribleViaKernelFn") {
runTestConstantGlobalVar();
runTestGlobalArray();
}
+248
Ver ficheiro
@@ -0,0 +1,248 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
#define LEN8 8 * 4
#define LEN9 9 * 4
#define LEN10 10 * 4
#define LEN11 11 * 4
#define LEN12 12 * 4
__global__ void MemCpy8(uint8_t* In, uint8_t* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memcpy(Out + tid * 8, In + tid * 8, 8);
}
__global__ void MemCpy9(uint8_t* In, uint8_t* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memcpy(Out + tid * 9, In + tid * 9, 9);
}
__global__ void MemCpy10(uint8_t* In, uint8_t* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memcpy(Out + tid * 10, In + tid * 10, 10);
}
__global__ void MemCpy11(uint8_t* In, uint8_t* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memcpy(Out + tid * 11, In + tid * 11, 11);
}
__global__ void MemCpy12(uint8_t* In, uint8_t* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memcpy(Out + tid * 12, In + tid * 12, 12);
}
__global__ void MemSet8(uint8_t* In) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memset(In + tid * 8, 1, 8);
}
__global__ void MemSet9(uint8_t* In) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memset(In + tid * 9, 1, 9);
}
__global__ void MemSet10(uint8_t* In) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memset(In + tid * 10, 1, 10);
}
__global__ void MemSet11(uint8_t* In) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memset(In + tid * 11, 1, 11);
}
__global__ void MemSet12(uint8_t* In) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
memset(In + tid * 12, 1, 12);
}
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check memcpy and memset via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/hipTestMemKernel.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_kernel_MemoryOperationsViaKernels") {
uint8_t *A, *Ad, *B, *Bd, *C, *Cd;
A = new uint8_t[LEN8];
B = new uint8_t[LEN8];
C = new uint8_t[LEN8];
for (uint32_t i = 0; i < LEN8; i++) {
A[i] = i;
B[i] = 0;
C[i] = 0;
}
HIP_CHECK(hipMalloc(&Ad, LEN8));
HIP_CHECK(hipMalloc(&Bd, LEN8));
HIP_CHECK(hipMalloc(&Cd, LEN8));
HIP_CHECK(hipMemcpy(Ad, A, LEN8, hipMemcpyHostToDevice));
hipLaunchKernelGGL(MemCpy8, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd);
hipLaunchKernelGGL(MemSet8, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd);
HIP_CHECK(hipMemcpy(B, Bd, LEN8, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, LEN8, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN8; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(C[i] == 1);
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
SECTION("MemCpySet1") {
A = new uint8_t[LEN9];
B = new uint8_t[LEN9];
C = new uint8_t[LEN9];
for (uint32_t i = 0; i < LEN9; i++) {
A[i] = i;
B[i] = 0;
C[i] = 0;
}
HIP_CHECK(hipMalloc(&Ad, LEN9));
HIP_CHECK(hipMalloc(&Bd, LEN9));
HIP_CHECK(hipMalloc(&Cd, LEN9));
HIP_CHECK(hipMemcpy(Ad, A, LEN9, hipMemcpyHostToDevice));
hipLaunchKernelGGL(MemCpy9, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd);
hipLaunchKernelGGL(MemSet9, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd);
HIP_CHECK(hipMemcpy(B, Bd, LEN9, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, LEN9, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN9; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(C[i] == 1);
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}
SECTION("MemCpySet2") {
A = new uint8_t[LEN10];
B = new uint8_t[LEN10];
C = new uint8_t[LEN10];
for (uint32_t i = 0; i < LEN10; i++) {
A[i] = i;
B[i] = 0;
C[i] = 0;
}
HIP_CHECK(hipMalloc(&Ad, LEN10));
HIP_CHECK(hipMalloc(&Bd, LEN10));
HIP_CHECK(hipMalloc(&Cd, LEN10));
HIP_CHECK(hipMemcpy(Ad, A, LEN10, hipMemcpyHostToDevice));
hipLaunchKernelGGL(MemCpy10, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd);
hipLaunchKernelGGL(MemSet10, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd);
HIP_CHECK(hipMemcpy(B, Bd, LEN10, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, LEN10, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN10; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(C[i] == 1);
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}
SECTION("MemCpySet3") {
A = new uint8_t[LEN11];
B = new uint8_t[LEN11];
C = new uint8_t[LEN11];
for (uint32_t i = 0; i < LEN11; i++) {
A[i] = i;
B[i] = 0;
C[i] = 0;
}
HIP_CHECK(hipMalloc(&Ad, LEN11));
HIP_CHECK(hipMalloc(&Bd, LEN11));
HIP_CHECK(hipMalloc(&Cd, LEN11));
HIP_CHECK(hipMemcpy(Ad, A, LEN11, hipMemcpyHostToDevice));
hipLaunchKernelGGL(MemCpy11, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd);
hipLaunchKernelGGL(MemSet11, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd);
HIP_CHECK(hipMemcpy(B, Bd, LEN11, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, LEN11, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN11; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(C[i] == 1);
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}
SECTION("MemCpySet4") {
A = new uint8_t[LEN12];
B = new uint8_t[LEN12];
C = new uint8_t[LEN12];
for (uint32_t i = 0; i < LEN12; i++) {
A[i] = i;
B[i] = 0;
C[i] = 0;
}
HIP_CHECK(hipMalloc(&Ad, LEN12));
HIP_CHECK(hipMalloc(&Bd, LEN12));
HIP_CHECK(hipMalloc(&Cd, LEN12));
HIP_CHECK(hipMemcpy(Ad, A, LEN12, hipMemcpyHostToDevice));
hipLaunchKernelGGL(MemCpy12, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd);
hipLaunchKernelGGL(MemSet12, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd);
HIP_CHECK(hipMemcpy(B, Bd, LEN12, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, LEN12, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN12; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(C[i] == 1);
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}
}
+136
Ver ficheiro
@@ -0,0 +1,136 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
#define NUM 1024
#define THREADS_PER_BLOCK_X 4
// Device (Kernel) function, it must be void
__global__ void vadd_asm(float* out, float* in) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
#ifdef __HIP_PLATFORM_NVIDIA__
asm volatile("add.f32 %0,%1,%2;" : "=f"(out[i]) : "f"(in[i]), "f"(out[i]));
#endif
#ifdef __HIP_PLATFORM_AMD__
asm volatile("v_add_f32_e32 %0, %1, %2" : "=v"(out[i]) : "v"(in[i]),
"v"(out[i]));
#endif
}
// CPU implementation of Vector Result
void addCPUReference(float* output, float* input) {
for (unsigned int j = 0; j < NUM; j++) {
output[j] = input[j] + output[j];
}
}
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check inline asm vadd instruction via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/inline_asm_vadd.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_kernel_inline_asm_vadd_Functional") {
float* VectorA;
float* ResultVector;
float* VectorB;
float* gpuVector;
float* gpuResultVector;
int i;
int errors;
VectorA = reinterpret_cast<float*>(malloc(NUM * sizeof(float)));
ResultVector = reinterpret_cast<float*>(malloc(NUM * sizeof(float)));
VectorB = reinterpret_cast<float*>(malloc(NUM * sizeof(float)));
// initialize the input data
for (i = 0; i < NUM; i++) {
VectorA[i] = static_cast<float>(i * 10.0f);
VectorB[i] = static_cast<float>(i * 30.0f);
}
// allocate the memory on the device side
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&gpuVector),
NUM * sizeof(float)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&gpuResultVector),
NUM * sizeof(float)));
// Memory transfer from host to device
HIP_CHECK(hipMemcpy(gpuVector, VectorA, NUM * sizeof(float),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(gpuResultVector, VectorB, NUM * sizeof(float),
hipMemcpyHostToDevice));
// Lauching kernel from host
hipLaunchKernelGGL(vadd_asm, dim3(NUM / THREADS_PER_BLOCK_X),
dim3(THREADS_PER_BLOCK_X), 0, 0,
gpuResultVector, gpuVector);
// Memory transfer from device to host
HIP_CHECK(hipMemcpy(ResultVector, gpuResultVector, NUM * sizeof(float),
hipMemcpyDeviceToHost));
// CPU Result computation
addCPUReference(VectorB, VectorA);
// verify the results
errors = 0;
double eps = 1.0E-3;
for (i = 0; i < NUM; i++) {
if (std::abs(ResultVector[i] - VectorB[i]) > eps) {
errors++;
}
}
if (errors != 0) {
REQUIRE(false);
} else {
REQUIRE(true);
}
// free the resources on device side
HIP_CHECK(hipFree(gpuVector));
HIP_CHECK(hipFree(gpuResultVector));
HIP_CHECK(hipDeviceReset());
// free the resources on host side
free(VectorA);
free(ResultVector);
free(VectorB);
}
+128
Ver ficheiro
@@ -0,0 +1,128 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
#define NUM 1024
#define THREADS_PER_BLOCK_X 4
// Device (Kernel) function, it must be void
__global__ void vmac_asm(float* out, float* in, float a) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
asm volatile("v_mac_f32_e32 %0, %2, %3" : "=v"(out[i]) :
"0"(out[i]), "v"(a), "v"(in[i]));
}
// CPU implementation of saxpy
void addCPUReference(float* output, float* input, float a) {
for (unsigned int j = 0; j < NUM; j++) {
output[j] = a * input[j] + output[j];
}
}
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check inline asm vmac instruction via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/inline_asm_vmac.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_kernel_inline_asm_vmac_Functional") {
float* VectorA;
float* ResultVector;
float* VectorB;
float* gpuVector;
float* gpuResultVector;
const float a = 10.0f;
int i;
int errors;
VectorA = reinterpret_cast<float*>(malloc(NUM * sizeof(float)));
ResultVector = reinterpret_cast<float*>(malloc(NUM * sizeof(float)));
VectorB = reinterpret_cast<float*>(malloc(NUM * sizeof(float)));
// initialize the input data
for (i = 0; i < NUM; i++) {
VectorA[i] = static_cast<float>(i * 10.0f);
VectorB[i] = static_cast<float>(i * 30.0f);
}
// allocate the memory on the device side
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&gpuVector),
NUM * sizeof(float)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&gpuResultVector),
NUM * sizeof(float)));
// Memory transfer from host to device
HIP_CHECK(hipMemcpy(gpuVector, VectorA, NUM * sizeof(float),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(gpuResultVector, VectorB, NUM * sizeof(float),
hipMemcpyHostToDevice));
// Lauching kernel from host
hipLaunchKernelGGL(vmac_asm, dim3(NUM / THREADS_PER_BLOCK_X),
dim3(THREADS_PER_BLOCK_X), 0, 0,
gpuResultVector, gpuVector, a);
// Memory transfer from device to host
HIP_CHECK(hipMemcpy(ResultVector, gpuResultVector, NUM * sizeof(float),
hipMemcpyDeviceToHost));
// CPU Result computation
addCPUReference(VectorB, VectorA, a);
// verify the results
errors = 0;
double eps = 1.0E-3;
for (i = 0; i < NUM; i++) {
if (std::abs(ResultVector[i] - VectorB[i]) > eps) {
errors++;
}
}
if (errors != 0) {
REQUIRE(false);
} else {
REQUIRE(true);
}
// free the resources on device side
HIP_CHECK(hipFree(gpuVector));
HIP_CHECK(hipFree(gpuResultVector));
HIP_CHECK(hipDeviceReset());
// free the resources on host side
free(VectorA);
free(ResultVector);
free(VectorB);
}
+102
Ver ficheiro
@@ -0,0 +1,102 @@
/*
Copyright (c) 2023 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, INCLUDING 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 ANY 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.
*/
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
constexpr size_t N = 1024;
int p_blockSize = 256;
__global__ void
__launch_bounds__(256, 2)
myKern(int* C, const int* A, int N) {
int tid = (blockIdx.x * blockDim.x + threadIdx.x);
if (tid < N) {
C[tid] = A[tid];
}
}
/**
* @addtogroup hipLaunchKernelGGL
* @{
* @ingroup KernelTest
* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` -
* Method to invocate kernel functions
*/
/**
* Test Description
* ------------------------
* - Test case to check launch bounds via kernel call.
* Test source
* ------------------------
* - catch/unit/kernel/launch_bounds.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_kernel_LaunchBounds_Functional") {
size_t Nbytes = N * sizeof(int);
int *A_d, *C_d, *A_h, *C_h;
HIPCHECK(hipMalloc(&A_d, Nbytes));
HIPCHECK(hipMalloc(&C_d, Nbytes));
A_h = reinterpret_cast<int*>(malloc(Nbytes));
C_h = reinterpret_cast<int*>(malloc(Nbytes));
for (int i = 0; i < N; i++) {
A_h[i] = i * 10;
C_h[i] = 0x0;
}
int blocks = N / p_blockSize;
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipGetLastError());
hipLaunchKernelGGL(myKern, dim3(blocks), dim3(p_blockSize), 0,
0, C_d, A_d, N);
#ifdef __HIP_PLATFORM_NVIDIA__
cudaFuncAttributes attrib;
cudaFuncGetAttributes(&attrib, myKern);
printf("binaryVersion = %d\n", attrib.binaryVersion);
printf("cacheModeCA = %d\n", attrib.cacheModeCA);
printf("constSizeBytes = %zu\n", attrib.constSizeBytes);
printf("localSizeBytes = %zud\n", attrib.localSizeBytes);
printf("maxThreadsPerBlock = %d\n", attrib.maxThreadsPerBlock);
printf("numRegs = %d\n", attrib.numRegs);
printf("ptxVersion = %d\n", attrib.ptxVersion);
printf("sharedSizeBytes = %zud\n", attrib.sharedSizeBytes);
#endif
HIPCHECK(hipDeviceSynchronize());
HIPCHECK(hipGetLastError());
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIPCHECK(hipDeviceSynchronize());
for (int i = 0; i < N; i++) {
int goldVal = i * 10;
REQUIRE(C_h[i] == goldVal);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
free(A_h);
free(C_h);
}
+99
Ver ficheiro
@@ -0,0 +1,99 @@
/*
Copyright (c) 2023 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.
*/
#ifndef _STRESSTEST_PRINTF_COMMON_H_
#define _STRESSTEST_PRINTF_COMMON_H_
#include <errno.h>
#include <error.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <math.h>
#include <fstream>
#include <iostream>
#include <string>
struct CaptureStream {
int saved_fd;
int orig_fd;
int temp_fd;
char tempname[13] = "mytestXXXXXX";
explicit CaptureStream(FILE *original) {
orig_fd = fileno(original);
saved_fd = dup(orig_fd);
if ((temp_fd = mkstemp(tempname)) == -1) {
error(0, errno, "Error");
assert(false);
}
fflush(nullptr);
if (dup2(temp_fd, orig_fd) == -1) {
error(0, errno, "Error");
assert(false);
}
if (close(temp_fd) != 0) {
error(0, errno, "Error");
assert(false);
}
}
void restoreStream() {
if (saved_fd == -1)
return;
fflush(nullptr);
if (dup2(saved_fd, orig_fd) == -1) {
error(0, errno, "Error");
assert(false);
}
if (close(saved_fd) != 0) {
error(0, errno, "Error");
assert(false);
}
saved_fd = -1;
}
const char *getTempFilename() {
return (const char*)tempname;
}
std::ifstream getCapturedData() {
restoreStream();
std::ifstream temp(tempname);
return temp;
}
~CaptureStream() {
restoreStream();
if (remove(tempname) != 0) {
error(0, errno, "Error");
assert(false);
}
}
};
#endif // _STRESSTEST_PRINTF_COMMON_H_