diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index de70022a06..2b086edcc6 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -42,6 +42,7 @@ add_subdirectory(module) add_subdirectory(channelDescriptor) add_subdirectory(executionControl) add_subdirectory(p2p) +add_subdirectory(gcc) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) diff --git a/projects/hip-tests/catch/unit/gcc/CMakeLists.txt b/projects/hip-tests/catch/unit/gcc/CMakeLists.txt new file mode 100644 index 0000000000..c7f7f8b1da --- /dev/null +++ b/projects/hip-tests/catch/unit/gcc/CMakeLists.txt @@ -0,0 +1,26 @@ +# Common Tests - Test independent of all platforms +if(HIP_PLATFORM MATCHES "amd") +if(UNIX) +set(TEST_SRC + gccTest.cc + gpu.cpp +) +# Creating Custom object file +add_custom_target(LaunchKernel_custom COMMAND gcc -c ${CMAKE_CURRENT_SOURCE_DIR}/LaunchKernel.c -I${HIP_PATH}/include -D__HIP_PLATFORM_AMD__ -o LaunchKernel.o) +add_custom_target(hipMalloc_custom COMMAND gcc -c ${CMAKE_CURRENT_SOURCE_DIR}/hipMalloc.c -I${HIP_PATH}/include -D__HIP_PLATFORM_AMD__ -o hipMalloc.o) + +add_library(LaunchKernel_lib OBJECT IMPORTED) +add_library(hipMalloc_lib OBJECT IMPORTED) + +set_property(TARGET LaunchKernel_lib PROPERTY IMPORTED_OBJECTS "${CMAKE_CURRENT_BINARY_DIR}/LaunchKernel.o") +set_property(TARGET hipMalloc_lib PROPERTY IMPORTED_OBJECTS "${CMAKE_CURRENT_BINARY_DIR}/hipMalloc.o") + + +hip_add_exe_to_target(NAME gccTests + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS LaunchKernel_lib hipMalloc_lib) + +add_dependencies(gccTests LaunchKernel_custom hipMalloc_custom) +endif() +endif() diff --git a/projects/hip-tests/catch/unit/gcc/LaunchKernel.c b/projects/hip-tests/catch/unit/gcc/LaunchKernel.c new file mode 100644 index 0000000000..beac8346ad --- /dev/null +++ b/projects/hip-tests/catch/unit/gcc/LaunchKernel.c @@ -0,0 +1,195 @@ + +/* Copyright (c) 2019 - 2021 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. + */ + +#include +#include +#include "LaunchKernel.h" +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" + +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError), \ + localError, #error, __FILE__, __LINE__, KNRM); \ + printf("API returned error code.\n"); \ + } \ + } + +bool LaunchKernelArg() +{ + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel), blocks, threads, NULL, 0, 0)); + + return true; +} + +bool LaunchKernelArg1() +{ + int A = 0; + int *A_d = NULL; + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + void* Args[]={&A_d}; + + HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel1), blocks, threads, Args, 0, 0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + + if(A != 333) + return false; + + return true; +} + +bool LaunchKernelArg2() +{ + int A = 0; + int B = 123; + int *A_d = NULL; + int *B_d = NULL; + + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&B_d, sizeof(int))); + + // Copy data from host memory to device memory + HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice)); + + void* Args[]={&A_d, &B_d}; + HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel2), blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(B_d)); + + if(A != 123) + return false; + + return true; +} + +bool LaunchKernelArg3() +{ + int A = 321; + int B = 123; + int C = 0; + int *A_d = NULL; + int *B_d = NULL; + int *C_d = NULL; + + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&B_d, sizeof(int))); + + HIPCHECK(hipMalloc((void**)&C_d, sizeof(int))); + + // Copy data from host memory to device memory + HIPCHECK(hipMemcpy(A_d, &A, sizeof(int), hipMemcpyHostToDevice)); + + HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice)); + + void* Args[]={&A_d, &B_d, &C_d}; + HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel3), blocks, threads, Args,0,0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(B_d)); + HIPCHECK(hipFree(C_d)); + + if(C != 444) + return false; + + return true; +} + +bool LaunchKernelArg4() +{ + int A = 0; + int *A_d = NULL; + dim3 blocks = {1,1,1}; + dim3 threads = {1,1,1}; + + // Allocate Device memory + HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); + + char c = 1; + short s = 10; + int i = 100; + struct things t = {2,20,200}; + + void* Args[]={&A_d, &c, &s, &i, &t}; + HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel4), blocks, threads, Args, 0, 0)); + + // Get the result back to host memory + HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + HIPCHECK(hipFree(A_d)); + + if (A != (c + s + i + t.c + t.s + t.i)) + return false; + + return true; +} + +#ifdef __cplusplus +extern "C" +{ +#endif +int launchKernel() +{ + printf("Calling LaunchKernel.c file\n"); + if( LaunchKernelArg() && + LaunchKernelArg1() && + LaunchKernelArg2() && + LaunchKernelArg3() && + LaunchKernelArg4()) + { + printf("PASSED!\n"); + return 1; + } + else + printf("FAILED\n"); + return 0; +} +#ifdef __cplusplus +} +#endif diff --git a/projects/hip-tests/catch/unit/gcc/LaunchKernel.h b/projects/hip-tests/catch/unit/gcc/LaunchKernel.h new file mode 100644 index 0000000000..e5f630924b --- /dev/null +++ b/projects/hip-tests/catch/unit/gcc/LaunchKernel.h @@ -0,0 +1,46 @@ + +/* Copyright (c) 2019 - 2021 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. + */ + +#ifdef __cplusplus +extern "C" +{ +#endif + +struct things { + char c; + short s; + int i; +}; + +typedef enum func{ + mykernel, + mykernel1, + mykernel2, + mykernel3, + mykernel4 +}func; + +extern const void* getKernelFunc(enum func f); + +int launchKernel(); +int hipMallocfunc(); + +#ifdef __cplusplus +} +#endif diff --git a/projects/hip-tests/catch/unit/gcc/gccTest.cc b/projects/hip-tests/catch/unit/gcc/gccTest.cc new file mode 100644 index 0000000000..6c64553558 --- /dev/null +++ b/projects/hip-tests/catch/unit/gcc/gccTest.cc @@ -0,0 +1,64 @@ +/* + * 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, 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. + * */ + +#include +#include +extern "C" { +#include "LaunchKernel.h" +} + +/** + * Test Description + * ------------------------ + * - calling launchKernel which is c function from catch2 + * and compile with gcc compiler and verify the results. + + * Test source + * ------------------------ + * - catch/unit/gcc/gccTest.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_LaunchKernelgccTests") { + printf("Calling launchKernel files from here\n"); + int result = launchKernel(); + REQUIRE(result == 1); +} + +/** + * Test Description + * ------------------------ + * - Calling hipMalloc which is c file from catch2 and compile + * with gcc compiler and verify the results. + + * Test source + * ------------------------ + * - catch/unit/gcc/gccTest.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipMallocgccTests") { + printf("Calling hipMalloc files from here\n"); + int result = hipMallocfunc(); + REQUIRE(result == 1); +} diff --git a/projects/hip-tests/catch/unit/gcc/gpu.cpp b/projects/hip-tests/catch/unit/gcc/gpu.cpp new file mode 100644 index 0000000000..e7c98e4117 --- /dev/null +++ b/projects/hip-tests/catch/unit/gcc/gpu.cpp @@ -0,0 +1,62 @@ +/* + * Copyright (c) 2019 - 2021 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. + * */ + +#include +#include +#include "LaunchKernel.h" + +extern "C" +{ + +__global__ void kernel() +{ +} + +__global__ void kernel1(int *a) +{ + *a = 333; +} + +__global__ void kernel2(int *a, int*b) +{ + *a = *b; +} + +__global__ void kernel3(int *a, int*b, int* c) +{ + *c = *a+*b; +} + +__global__ void kernel4(int *a, char c, short s, int i, struct things t) +{ + *a = c + s + i + t.c + t.s + t.i; +} + +const void* funcTable[] = { + (const void*)kernel, + (const void*)kernel1, + (const void*)kernel2, + (const void*)kernel3, + (const void*)kernel4 }; + +const void* getKernelFunc(enum func f){ + return funcTable[f]; + } + +}//extern "C" diff --git a/projects/hip-tests/catch/unit/gcc/hipMalloc.c b/projects/hip-tests/catch/unit/gcc/hipMalloc.c new file mode 100644 index 0000000000..5b6ad0cd86 --- /dev/null +++ b/projects/hip-tests/catch/unit/gcc/hipMalloc.c @@ -0,0 +1,35 @@ +/* + * Copyright (c) 2019 - 2021 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. + * */ + +#include +#include + +#ifdef __cplusplus +extern "C" +{ +#endif +int hipMallocfunc() { + int *Ad; + hipMalloc((void**)&Ad, 1024); + printf("hipMalloc PASSED!\n"); + return 1; +} +#ifdef __cplusplus +} +#endif