diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 31bd37da9..da2d22894 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -35,6 +35,7 @@ add_subdirectory(compiler) add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) add_subdirectory(context) +add_subdirectory(module) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) #add_subdirectory(clock) diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt new file mode 100644 index 000000000..eda33f8ef --- /dev/null +++ b/catch/unit/module/CMakeLists.txt @@ -0,0 +1,43 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipModuleNegative.cc + hipModuleGetGlobal.cc + hipModuleOccupancyMaxPotentialActiveBlockSize.cc + hipModuleLoadMultiThreaded.cc + hipModuleLoadDataMultThreaded.cc + hipModuleLoadDataMultThreadOnMultGPU.cc + hipModuleLoadUnloadStress.cc + hipModuleTexture2dDrv.cc +) + +add_custom_target(vcpy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/vcpy_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/vcpy_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +add_custom_target(global_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/global_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/global_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +add_custom_target(empty_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/empty_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/empty_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +add_custom_target(kernel_composite_test.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_composite_test.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_composite_test.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +add_custom_target(tex2d_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/tex2d_kernel.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../module/tex2d_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +if (UNIX) +add_custom_target(opencl_add.cc COMMAND cp ${CMAKE_CURRENT_SOURCE_DIR}/opencl_add.cc ${CMAKE_CURRENT_BINARY_DIR}) +endif() + +# only for AMD +if(HIP_PLATFORM MATCHES "amd") + set(AMD_SRC + hipOpenCLCOTest.cc + ) + set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) +endif() + +hip_add_exe_to_target(NAME ModuleTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) + +if (UNIX) +add_dependencies(build_tests vcpy_kernel.code global_kernel.code empty_kernel.code kernel_composite_test.code tex2d_kernel.code opencl_add.cc) +else() +add_dependencies(build_tests vcpy_kernel.code global_kernel.code empty_kernel.code kernel_composite_test.code tex2d_kernel.code) +endif() diff --git a/catch/unit/module/empty_kernel.cc b/catch/unit/module/empty_kernel.cc new file mode 100644 index 000000000..252930556 --- /dev/null +++ b/catch/unit/module/empty_kernel.cc @@ -0,0 +1,24 @@ +/* +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. +*/ + +#include "hip/hip_runtime.h" + +extern "C" __global__ void EmptyKernel() { +} + diff --git a/catch/unit/module/global_kernel.cc b/catch/unit/module/global_kernel.cc new file mode 100644 index 000000000..82db73fa9 --- /dev/null +++ b/catch/unit/module/global_kernel.cc @@ -0,0 +1,36 @@ +/* +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. +*/ + +#include "hip/hip_runtime.h" + +#define ARRAY_SIZE (16) + +__device__ float myDeviceGlobal; +__device__ float myDeviceGlobalArray[16]; + + +extern "C" __global__ void hello_world(const float* a, float* b) { + int tx = threadIdx.x; + b[tx] = a[tx]; +} + +extern "C" __global__ void test_globals(const float* a, float* b) { + int tx = threadIdx.x; + b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE]; +} diff --git a/catch/unit/module/hipModuleGetGlobal.cc b/catch/unit/module/hipModuleGetGlobal.cc new file mode 100644 index 000000000..576bfd347 --- /dev/null +++ b/catch/unit/module/hipModuleGetGlobal.cc @@ -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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#include +#include + +#define fileName "global_kernel.code" +#define LEN 64 +#define SIZE LEN * sizeof(float) +#define ARRAY_SIZE 16 + +struct { + void* _Ad; + void* _Bd; +} args; + +TEST_CASE("Unit_hipModuleGetGlobal") { + float *A, *B; + float *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + #if HT_NVIDIA + HIP_CHECK(hipCtxCreate(&context, 0, device)); + #endif + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(Ad), A, SIZE)); + HIP_CHECK(hipMemcpyHtoD((hipDeviceptr_t)(Bd), B, SIZE)); + hipModule_t Module; + HIP_CHECK(hipModuleLoad(&Module, fileName)); + + float myDeviceGlobal_h = 42.0; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; + HIP_CHECK(hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, + Module, "myDeviceGlobal")); + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, + deviceGlobalSize)); + + float myDeviceGlobalArray_h[ARRAY_SIZE]; + hipDeviceptr_t myDeviceGlobalArray; + size_t myDeviceGlobalArraySize; + + HIP_CHECK(hipModuleGetGlobal(reinterpret_cast(&myDeviceGlobalArray), //NOLINT + &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray")); + + for (int i = 0; i < ARRAY_SIZE; i++) { + myDeviceGlobalArray_h[i] = i * 1000.0f; + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), + &myDeviceGlobalArray_h, myDeviceGlobalArraySize)); + } + + args._Ad = reinterpret_cast(Ad); + args._Bd = reinterpret_cast(Bd); + + size_t size = sizeof(args); + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; + + SECTION("Running test for hello world kernel") { + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world")); + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, + reinterpret_cast(&config))); + + HIP_CHECK(hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE)); + + int mismatchCount = 0; + for (uint32_t i = 0; i < LEN; i++) { + if (A[i] != B[i]) { + mismatchCount++; + if (mismatchCount >= 10) { + break; + } + } + } + REQUIRE(mismatchCount == 0); + } + + SECTION("running test for tests_globals kernel") { + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, + "test_globals")); + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, + reinterpret_cast(&config))); + HIP_CHECK(hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE)); + + int mismatchCount = 0; + for (uint32_t i = 0; i < LEN; i++) { + float expected; + expected = A[i] + myDeviceGlobal_h + + myDeviceGlobalArray_h[i % 16]; + if (expected != B[i]) { + mismatchCount++; + if (mismatchCount >= 10) { + break; + } + } + } + REQUIRE(mismatchCount == 0); + } + + HIP_CHECK(hipModuleUnload(Module)); + #if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); + #endif +} diff --git a/catch/unit/module/hipModuleLoadDataMultThreadOnMultGPU.cc b/catch/unit/module/hipModuleLoadDataMultThreadOnMultGPU.cc new file mode 100644 index 000000000..c1d45f575 --- /dev/null +++ b/catch/unit/module/hipModuleLoadDataMultThreadOnMultGPU.cc @@ -0,0 +1,142 @@ +/* +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 +#include +#include +#include +#include + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 8 + +#define FILENAME "vcpy_kernel.code" +#define kernel_name "hello_world" + +static std::vector load_file() { + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + printf("Info:could not open code object '%s'\n", FILENAME); + } + return buffer; +} + +static void run(const std::vector& buffer, int device) { + HIP_CHECK(hipSetDevice(device)); + hipModule_t Module; + hipFunction_t Function; + + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + HIP_CHECK(hipMalloc(&Ad, SIZE)); + HIP_CHECK(hipMalloc(&Bd, SIZE)); + + HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + + HIP_CHECK(hipModuleLoadData(&Module, &buffer[0])); + HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = static_cast(Ad); + args._Bd = static_cast(Bd); + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, + NULL, reinterpret_cast(&config))); + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipModuleUnload(Module)); + HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + REQUIRE(A[i] == B[i]); + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + delete[] A; + delete[] B; +} + +struct joinable_thread : std::thread { + template + joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) {} // NOLINT + + joinable_thread& operator=(joinable_thread&& other) = default; + joinable_thread(joinable_thread&& other) = default; + + ~joinable_thread() { + if (this->joinable()) + this->join(); + } +}; + +static void run_multi_threads(uint32_t n, const std::vector& buffer) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + std::vector threads; + + for (int device =0; device < numDevices; ++device) { + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[&, device] { + run(buffer, device); + } + }); + } + } +} + +TEST_CASE("Unit_hipModuleLoadDataMultGPUOnMultThread") { + HIP_CHECK(hipInit(0)); + auto buffer = load_file(); + auto file_size = buffer.size() / (1024 * 1024); + auto thread_count = HipTest::getHostThreadCount(file_size + 10, THREADS); + if (thread_count == 0) { + HipTest::HIP_SKIP_TEST("Skipping because thread_count is 0"); + return; + } + // run multi thread on multi devices + run_multi_threads(thread_count, buffer); +} diff --git a/catch/unit/module/hipModuleLoadDataMultThreaded.cc b/catch/unit/module/hipModuleLoadDataMultThreaded.cc new file mode 100644 index 000000000..268c7bc9f --- /dev/null +++ b/catch/unit/module/hipModuleLoadDataMultThreaded.cc @@ -0,0 +1,135 @@ +/* +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 +#include +#include +#include +#include + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 8 + +#define FILENAME "vcpy_kernel.code" +#define kernel_name "hello_world" + +static std::vector load_file() { + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + printf("Info:could not open code object '%s'\n", FILENAME); + } + return buffer; +} + +void run(const std::vector& buffer) { + hipModule_t Module; + hipFunction_t Function; + + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + + HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + + HIP_CHECK(hipModuleLoadData(&Module, &buffer[0])); + HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = reinterpret_cast(Ad); + args._Bd = reinterpret_cast(Bd); + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, + NULL, reinterpret_cast(&config))); + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipModuleUnload(Module)); + HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + REQUIRE(A[i] == B[i]); + } + + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + delete[] A; + delete[] B; +} + +struct joinable_thread : std::thread { + template + joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) {} // NOLINT + + joinable_thread& operator=(joinable_thread&& other) = default; + joinable_thread(joinable_thread&& other) = default; + + ~joinable_thread() { + if (this->joinable()) + this->join(); + } +}; + +static void run_multi_threads(uint32_t n, const std::vector& buffer) { + std::vector threads; + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[&] { + run(buffer); + } + }); + } +} + +TEST_CASE("Unit_hipModuleLoadDataMultThreaded") { + HIP_CHECK(hipInit(0)); + auto buffer = load_file(); + auto file_size = buffer.size() / (1024 * 1024); + auto thread_count = HipTest::getHostThreadCount(file_size + 10, THREADS); + if (thread_count == 0) { + HipTest::HIP_SKIP_TEST("Skipping because thread_count is 0"); + return; + } + // run multi threads + run_multi_threads(thread_count, buffer); +} diff --git a/catch/unit/module/hipModuleLoadMultiThreaded.cc b/catch/unit/module/hipModuleLoadMultiThreaded.cc new file mode 100644 index 000000000..ea257ac30 --- /dev/null +++ b/catch/unit/module/hipModuleLoadMultiThreaded.cc @@ -0,0 +1,95 @@ +/* +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. +*/ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#define THREADS 8 +#define MAX_NUM_THREADS 512 + +#define WARMUP_RUN_COUNT 10 +#define TIMING_RUN_COUNT 100 +#define TOTAL_RUN_COUNT WARMUP_RUN_COUNT + TIMING_RUN_COUNT +#define FILENAME "empty_kernel.code" +#define kernel_name "EmptyKernel" + +void hipModuleLaunchKernel_enqueue(std::atomic_int* shared, int max_threads) { + // resources necessary for this thread + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoad(&module, FILENAME)); + HIP_CHECK(hipModuleGetFunction(&function, module, kernel_name)); + void* kernel_params = nullptr; + while (max_threads != shared->load(std::memory_order_acquire)) { + break; + } + + for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, stream, + &kernel_params, nullptr)); + } + HIP_CHECK(hipModuleUnload(module)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +// thread pool +struct thread_pool { + thread_pool(int total_threads) : max_threads(total_threads) { + } + void start(std::function f) { + for (int i = 0; i < max_threads; ++i) { + threads.push_back(std::async(std::launch::async, f, &shared, + max_threads)); + } + } + void finish() { + for (auto&&thread : threads) { + thread.get(); + } + threads.clear(); + shared = 0; + } + ~thread_pool() { + finish(); + } + private: + std::atomic_int shared {0}; + std::vector buffer; + std::vector> threads; + int max_threads = 1; +}; + +TEST_CASE("Unit_hipModuleLoadMultiThreaded") { + int max_threads = min(THREADS * std::thread::hardware_concurrency(), + MAX_NUM_THREADS); + thread_pool task(max_threads); + + task.start(hipModuleLaunchKernel_enqueue); + task.finish(); +} diff --git a/catch/unit/module/hipModuleLoadUnloadStress.cc b/catch/unit/module/hipModuleLoadUnloadStress.cc new file mode 100644 index 000000000..f3ee541d8 --- /dev/null +++ b/catch/unit/module/hipModuleLoadUnloadStress.cc @@ -0,0 +1,104 @@ +/* +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. +*/ + +#include +#include +#include +#include +#include +#include + +#define TEST_ITERATIONS 1000 +#define CODEOBJ_FILE "kernel_composite_test.code" + +/** + * Internal Function + */ +static std::vector load_file() { + std::ifstream file(CODEOBJ_FILE, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + printf("Info:could not open code object '%s'\n", CODEOBJ_FILE); + } + file.close(); + return buffer; +} +/** + * Validates no memory leakage for hipModuleLoad + */ +static void testhipModuleLoadUnloadStress() { + for (int count = 0; count < TEST_ITERATIONS; count++) { + hipModule_t Module; + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy")); + HIP_CHECK(hipModuleUnload(Module)); + } +} +/** + * Validates no memory leakage for hipModuleLoadData + */ +static void testhipModuleLoadDataUnloadStress() { + auto buffer = load_file(); + for (int count = 0; count < TEST_ITERATIONS; count++) { + hipModule_t Module; + HIP_CHECK(hipModuleLoadData(&Module, &buffer[0])); + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy")); + HIP_CHECK(hipModuleUnload(Module)); + } +} +/** + * Validates no memory leakage for hipModuleLoadDataEx + */ +static void testhipModuleLoadDataExUnloadStress() { + auto buffer = load_file(); + for (int count = 0; count < TEST_ITERATIONS; count++) { + hipModule_t Module; + HIP_CHECK(hipModuleLoadDataEx(&Module, &buffer[0], 0, + nullptr, nullptr)); + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy")); + HIP_CHECK(hipModuleUnload(Module)); + } +} + +TEST_CASE("Unit_hipModuleLoadUnloadStress") { +#if HT_NVIDIA + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); +#endif + SECTION("running hipModuleLoadUnloadStress") { + testhipModuleLoadUnloadStress(); + } + SECTION("running hipModuleLoadDataUnloadStress") { + testhipModuleLoadDataUnloadStress(); + } + SECTION("running hipModuleLoadDataExUnloadStress") { + testhipModuleLoadDataExUnloadStress(); + } +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif +} diff --git a/catch/unit/module/hipModuleNegative.cc b/catch/unit/module/hipModuleNegative.cc new file mode 100644 index 000000000..814c84cc4 --- /dev/null +++ b/catch/unit/module/hipModuleNegative.cc @@ -0,0 +1,770 @@ +/* +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. +*/ + +#include +#include +#include +#include +#include +#include +#include +#include + + +#define FILENAME_NONEXST "sample_nonexst.code" +#define FILENAME_EMPTY "emptyfile.code" +#define FILENAME_RAND "rand_file.code" +#define RANDOMFILE_LEN 2048 +#define CODEOBJ_FILE "vcpy_kernel.code" +#define KERNEL_NAME "hello_world" +#define KERNEL_NAME_NONEXST "xyz" +#define CODEOBJ_GLOBAL "global_kernel.code" +#define DEVGLOB_VAR_NONEXIST "xyz" +#define DEVGLOB_VAR "myDeviceGlobal" +/** + * Internal Function + */ +static std::vector load_file(const char* filename) { + std::ifstream file(filename, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + printf("Info:could not open code object '%s'\n", filename); + } + file.close(); + return buffer; +} + +/** + * Internal Function + */ +void createRandomFile(const char* filename) { + std::ofstream outfile(filename, std::ios::binary); + char buf[RANDOMFILE_LEN]; + unsigned int seed = 1; + for (int i = 0; i < RANDOMFILE_LEN; i++) { + buf[i] = HipTest::RAND_R(&seed) % 256; + } + outfile.write(buf, RANDOMFILE_LEN); + outfile.close(); +} + +/** + * Validates negative scenarios for hipModuleLoad + * module = nullptr + */ +bool testhipModuleLoadNeg1() { + bool TestPassed = false; + hipError_t ret; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoad(nullptr, CODEOBJ_FILE); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoad + * fname = nullptr + */ +bool testhipModuleLoadNeg2() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoad(&Module, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} +/** + * Validates negative scenarios for hipModuleLoad + * fname = empty file + */ +bool testhipModuleLoadNeg3() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + // Create an empty + std::fstream fs; + fs.open(FILENAME_EMPTY, std::ios::out); + fs.close(); +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoad(&Module, FILENAME_EMPTY); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + remove(FILENAME_EMPTY); + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoad + * fname = ramdom file + */ +bool testhipModuleLoadNeg4() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + // Create a binary file with random numbers + createRandomFile(FILENAME_RAND); +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoad(&Module, FILENAME_RAND); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + remove(FILENAME_RAND); + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoad + * fname = non existent file + */ +bool testhipModuleLoadNeg5() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoad(&Module, FILENAME_NONEXST); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoad + * fname = empty string "" + */ +bool testhipModuleLoadNeg6() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoad(&Module, ""); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoadData + * module = nullptr + */ +bool testhipModuleLoadDataNeg1() { + bool TestPassed = false; + hipError_t ret; + auto buffer = load_file(CODEOBJ_FILE); +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoadData(nullptr, &buffer[0]); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoadData + * image = nullptr + */ +bool testhipModuleLoadDataNeg2() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoadData(&Module, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoadData + * image = ramdom file + */ +bool testhipModuleLoadDataNeg3() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + // Create a binary file with random numbers + createRandomFile(FILENAME_RAND); + // Open the code object file and copy it in a buffer + auto buffer = load_file(FILENAME_RAND); +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoadData(&Module, &buffer[0]); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + remove(FILENAME_RAND); + return TestPassed; +} +/** + * Validates negative scenarios for hipModuleLoadDataEx + * module = nullptr + */ +bool testhipModuleLoadDataExNeg1() { + bool TestPassed = false; + hipError_t ret; + // Open the code object file and copy it in a buffer + auto buffer = load_file(CODEOBJ_FILE); +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoadDataEx(nullptr, &buffer[0], 0, nullptr, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoadDataEx + * image = nullptr + */ +bool testhipModuleLoadDataExNeg2() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoadDataEx(&Module, nullptr, 0, nullptr, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleLoadDataEx + * image = ramdom file + */ +bool testhipModuleLoadDataExNeg3() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + // Create a binary file with random numbers + createRandomFile(FILENAME_RAND); + // Open the code object file and copy it in a buffer + auto buffer = load_file(FILENAME_RAND); +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleLoadDataEx(&Module, &buffer[0], 0, nullptr, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + remove(FILENAME_RAND); + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetFunction + * Function = nullptr + */ +bool testhipModuleGetFunctionNeg1() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + ret = hipModuleGetFunction(nullptr, Module, KERNEL_NAME); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetFunction + * Module is uninitialized + */ +bool testhipModuleGetFunctionNeg2() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module = nullptr; + hipFunction_t Function; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetFunction + * kname = non existing function + */ +bool testhipModuleGetFunctionNeg3() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipFunction_t Function; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME_NONEXST); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetFunction + * kname = nullptr + */ +bool testhipModuleGetFunctionNeg4() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipFunction_t Function; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + ret = hipModuleGetFunction(&Function, Module, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetFunction + * Module = Unloaded Module + */ +bool testhipModuleGetFunctionNeg5() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipFunction_t Function; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + HIP_CHECK(hipModuleUnload(Module)); + ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetFunction + * kname = Empty String "" + */ +bool testhipModuleGetFunctionNeg6() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipFunction_t Function; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + ret = hipModuleGetFunction(&Function, Module, ""); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * dptr = nullptr + */ +bool testhipModuleGetGlobalNeg1() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + size_t deviceGlobalSize; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL)); + ret = hipModuleGetGlobal(nullptr, &deviceGlobalSize, Module, DEVGLOB_VAR); + REQUIRE(ret == hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * bytes = nullptr + */ +bool testhipModuleGetGlobalNeg2() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipDeviceptr_t deviceGlobal; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL)); + ret = hipModuleGetGlobal(&deviceGlobal, nullptr, Module, DEVGLOB_VAR); + REQUIRE(ret == hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * name = nullptr + */ +bool testhipModuleGetGlobalNeg3() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL)); + ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, nullptr); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * name = wrong name + */ +bool testhipModuleGetGlobalNeg4() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL)); + ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, + DEVGLOB_VAR_NONEXIST); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * name = Empty String "" + */ +bool testhipModuleGetGlobalNeg5() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL)); + ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, ""); + REQUIRE(ret != hipSuccess); + TestPassed = true; + HIP_CHECK(hipModuleUnload(Module)); +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * Module = Unloaded Module + */ +bool testhipModuleGetGlobalNeg6() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL)); + HIP_CHECK(hipModuleUnload(Module)); + ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, + DEVGLOB_VAR); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetGlobal + * Module = Uninitialized Module + */ +bool testhipModuleGetGlobalNeg7() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module = nullptr; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, + Module, DEVGLOB_VAR); + REQUIRE(ret != hipSuccess); + TestPassed = true; +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleUnload + * 1. Unload an uninitialized module + * 2. Unload an unloaded module + */ +bool testhipModuleLoadNeg7() { + bool TestPassed = false; + hipError_t ret; + hipModule_t Module = nullptr; +#if HT_NVIDIA + hipCtx_t context; + initHipCtx(&context); +#endif + // test case 1 + SECTION("No obj file") { + ret = hipModuleUnload(Module); + REQUIRE(ret != hipSuccess); + TestPassed = true; + } + // test case 2 + SECTION("CODEOBJ file") { + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + HIP_CHECK(hipModuleUnload(Module)); + ret = hipModuleUnload(Module); + REQUIRE(ret != hipSuccess); + TestPassed = true; + } +#if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); +#endif + return TestPassed; +} + +TEST_CASE("Unit_hipModuleNegative") { + bool TestPassed = true; + + SECTION("test running for testhipModuleLoadNeg1") { + REQUIRE(TestPassed == testhipModuleLoadNeg1()); + } + SECTION("test running for testhipModuleLoadNeg2") { + REQUIRE(TestPassed == testhipModuleLoadNeg2()); + } + SECTION("test running for testhipModuleLoadNeg3") { + REQUIRE(TestPassed == testhipModuleLoadNeg3()); + } + SECTION("test running for testhipModuleLoadNeg4") { + REQUIRE(TestPassed == testhipModuleLoadNeg4()); + } + SECTION("test running for testhipModuleLoadNeg5") { + REQUIRE(TestPassed == testhipModuleLoadNeg5()); + } + SECTION("test running for testhipModuleLoadNeg6") { + REQUIRE(TestPassed == testhipModuleLoadNeg6()); + } + SECTION("test running for testhipModuleLoadDataNeg1") { + REQUIRE(TestPassed == testhipModuleLoadDataNeg1()); + } + SECTION("test running for testhipModuleLoadDataNeg2") { + REQUIRE(TestPassed == testhipModuleLoadDataNeg2()); + } + SECTION("test running for testhipModuleLoadDataNeg3") { + REQUIRE(TestPassed == testhipModuleLoadDataNeg3()); + } + SECTION("test running for testhipModuleLoadDataExNeg1") { + REQUIRE(TestPassed == testhipModuleLoadDataExNeg1()); + } + SECTION("test running for testhipModuleLoadDataExNeg2") { + REQUIRE(TestPassed == testhipModuleLoadDataExNeg2()); + } + SECTION("test running for testhipModuleLoadDataExNeg3") { + REQUIRE(TestPassed == testhipModuleLoadDataExNeg3()); + } + SECTION("test running for testhipModuleGetFunctionNeg1") { + REQUIRE(TestPassed == testhipModuleGetFunctionNeg1()); + } + SECTION("test running for testhipModuleGetFunctionNeg2") { + REQUIRE(TestPassed == testhipModuleGetFunctionNeg2()); + } + SECTION("test running for testhipModuleGetFunctionNeg3") { + REQUIRE(TestPassed == testhipModuleGetFunctionNeg3()); + } + SECTION("test running for testhipModuleGetFunctionNeg4") { + REQUIRE(TestPassed == testhipModuleGetFunctionNeg4()); + } + #if HT_AMD + SECTION("test running for testhipModuleGetFunctionNeg5") { + REQUIRE(TestPassed == testhipModuleGetFunctionNeg5()); + } + #endif + SECTION("test running for testhipModuleGetFunctionNeg6") { + REQUIRE(TestPassed == testhipModuleGetFunctionNeg6()); + } + SECTION("test running for testhipModuleGetGlobalNeg1") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg1()); + } + SECTION("test running for testhipModuleGetGlobalNeg2") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg2()); + } + SECTION("test running for testhipModuleGetGlobalNeg3") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg3()); + } + SECTION("test running for testhipModuleGetGlobalNeg4") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg4()); + } + SECTION("test running for testhipModuleGetGlobalNeg5") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg5()); + } + #if HT_AMD + SECTION("test running for testhipModuleGetGlobalNeg6") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg6()); + } + SECTION("test running for testhipModuleGetGlobalNeg7") { + REQUIRE(TestPassed == testhipModuleGetGlobalNeg7()); + } + SECTION("test running for testhipModuleLoadNeg7") { + REQUIRE(TestPassed == testhipModuleLoadNeg7()); + } + #endif +} diff --git a/catch/unit/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cc b/catch/unit/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cc new file mode 100644 index 000000000..e6f9e73ae --- /dev/null +++ b/catch/unit/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cc @@ -0,0 +1,53 @@ +/* +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. +*/ + +#include +#include +#include + +#define fileName "vcpy_kernel.code" +#define kernel_name "hello_world" + +TEST_CASE("Unit_hipModuleOccupancyMaxPotentialActiveBlockSize") { + int gridSize = 0; + int blockSize = 0; + int numBlock = 0; + HIP_CHECK(hipInit(0)); + + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + #if HT_NVIDIA + HIP_CHECK(hipCtxCreate(&context, 0, device)); + #endif + hipModule_t Module; + hipFunction_t Function; + HIP_CHECK(hipModuleLoad(&Module, fileName)); + HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, + Function, 0, 0)); + assert(gridSize != 0 && blockSize != 0); + HIP_CHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, + Function, blockSize, 0)); + assert(numBlock != 0); + HIP_CHECK(hipModuleUnload(Module)); + #if HT_NVIDIA + HIP_CHECK(hipCtxDestroy(context)); + #endif +} diff --git a/catch/unit/module/hipModuleTexture2dDrv.cc b/catch/unit/module/hipModuleTexture2dDrv.cc new file mode 100644 index 000000000..7b0d85680 --- /dev/null +++ b/catch/unit/module/hipModuleTexture2dDrv.cc @@ -0,0 +1,583 @@ +/* +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. +*/ + + +#include +#include +#include + +#include +#include +#include +#include +#include + + +#define CODEOBJ_FILE "tex2d_kernel.code" +#define NON_EXISTING_TEX_NAME "xyz" +#define EMPTY_TEX_NAME "" +#define GLOBAL_KERNEL_VAR "deviceGlobalFloat" +#define TEX_REF "ftex" +#define WIDTH 256 +#define HEIGHT 256 +#define MAX_STREAMS 4 +#define GRIDDIMX 16 +#define GRIDDIMY 16 +#define GRIDDIMZ 1 +#define BLOCKDIMZ 1 + +#if HT_NVIDIA + +#define CTX_CREATE() \ + hipCtx_t context;\ + initHipCtx(&context); + +#define CTX_DESTROY() HIP_CHECK(hipCtxDestroy(context)); +#define HIP_TEX_REFERENCE hipTexRef +#define HIP_ARRAY hiparray +#else +#define CTX_CREATE() +#define CTX_DESTROY() +#define HIP_TEX_REFERENCE textureReference* +#define HIP_ARRAY hipArray* +#endif + +std::atomic g_thTestPassed(1); + +/** + * Validates negative scenarios for hipModuleGetTexRef + * texRef = nullptr + */ +bool testTexRefEqNullPtr() { + bool TestPassed = false; + hipModule_t Module; + CTX_CREATE() + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + REQUIRE(hipSuccess != hipModuleGetTexRef(nullptr, Module, "tex")); + TestPassed = true; + CTX_DESTROY() + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetTexRef + * name = nullptr + */ +bool testNameEqNullPtr() { + bool TestPassed = false; + hipModule_t Module; + HIP_TEX_REFERENCE texref; + CTX_CREATE() + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + REQUIRE(hipSuccess != hipModuleGetTexRef(&texref, Module, nullptr)); + TestPassed = true; + CTX_DESTROY() + return TestPassed; +} +/** + * Validates negative scenarios for hipModuleGetTexRef + * name = Non Existing Tex Name + */ +bool testInvalidTexName() { + bool TestPassed = false; + hipModule_t Module; + HIP_TEX_REFERENCE texref; + CTX_CREATE() + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + REQUIRE(hipSuccess != hipModuleGetTexRef(&texref, Module, + NON_EXISTING_TEX_NAME)); + TestPassed = true; + CTX_DESTROY() + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetTexRef + * name = Empty Tex Name + */ +bool testEmptyTexName() { + bool TestPassed = false; + hipModule_t Module; + HIP_TEX_REFERENCE texref; + CTX_CREATE() + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + REQUIRE(hipSuccess != hipModuleGetTexRef(&texref, Module, EMPTY_TEX_NAME)); + TestPassed = true; + CTX_DESTROY() + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetTexRef + * name = Global Kernel Variable + */ +bool testWrongTexRef() { + bool TestPassed = false; + hipModule_t Module; + HIP_TEX_REFERENCE texref; + CTX_CREATE() + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + REQUIRE(hipSuccess != hipModuleGetTexRef(&texref, Module, GLOBAL_KERNEL_VAR)); + TestPassed = true; + CTX_DESTROY() + return TestPassed; +} + +/** + * Validates negative scenarios for hipModuleGetTexRef + * module = unloaded module + */ +bool testUnloadedMod() { + bool TestPassed = false; + hipModule_t Module; + HIP_TEX_REFERENCE texref; + CTX_CREATE() + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + HIP_CHECK(hipModuleUnload(Module)); + REQUIRE(hipSuccess != hipModuleGetTexRef(&texref, Module, TEX_REF)); + TestPassed = true; + CTX_DESTROY() + return TestPassed; +} +/** + * Internal Functions + * + */ +std::vector load_file() { + std::ifstream file(CODEOBJ_FILE, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + printf("Info:could not open code object '%s'\n", CODEOBJ_FILE); + } + return buffer; +} + +template void fillTestBuffer(unsigned int width, + unsigned int height, + T* hData) { + if (std::is_same::value) { + for (unsigned int i = 0; i < height; i++) { + for (unsigned int j = 0; j < width; j++) { + hData[i * width + j] = i * width + j + 0.5; + } + } + } else if (std::is_same::value) { + for (unsigned int i = 0; i < height; i++) { + for (unsigned int j = 0; j < width; j++) { + hData[i * width + j] = i * width + j; + } + } + } else if (std::is_same::value) { // cpplint asking to make int16 instead of short. + for (unsigned int i = 0; i < height; i++) { + for (unsigned int j = 0; j < width; j++) { + hData[i * width + j] = (i * width + j)% + (std::numeric_limits::max()); + } + } + } else if (std::is_same::value) { + for (unsigned int i = 0; i < height; i++) { + for (unsigned int j = 0; j < width; j++) { + hData[i * width + j] = (i * width + j)% + (std::numeric_limits::max()); + } + } + } +} + +void allocInitArray(unsigned int width, + unsigned int height, + hipArray_Format format, + HIP_ARRAY* array + ) { + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = format; + desc.NumChannels = 1; + desc.Width = width; + desc.Height = height; + HIP_CHECK(hipArrayCreate(array, &desc)); +} + +template void copyBuffer2Array(unsigned int width, + unsigned int height, + T* hData, + T1 array + ) { + hip_Memcpy2D copyParam; + memset(©Param, 0, sizeof(copyParam)); +#if HT_NVIDIA + copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY; + copyParam.srcMemoryType = CU_MEMORYTYPE_HOST; + copyParam.dstArray = *array; +#else + copyParam.dstMemoryType = hipMemoryTypeArray; + copyParam.srcMemoryType = hipMemoryTypeHost; + copyParam.dstArray = array; +#endif + copyParam.srcHost = hData; + copyParam.srcPitch = width * sizeof(T); + copyParam.WidthInBytes = copyParam.srcPitch; + copyParam.Height = height; + HIP_CHECK(hipMemcpyParam2D(©Param)); +} + +template void assignArray2TexRef(hipArray_Format format, + const char* texRefName, + hipModule_t Module, + T array + ) { + HIP_TEX_REFERENCE texref; +#if HT_NVIDIA + HIP_CHECK(hipModuleGetTexRef(&texref, Module, texRefName)); + HIP_CHECK(hipTexRefSetAddressMode(texref, 0, CU_TR_ADDRESS_MODE_WRAP)); + HIP_CHECK(hipTexRefSetAddressMode(texref, 1, CU_TR_ADDRESS_MODE_WRAP)); + HIP_CHECK(hipTexRefSetFilterMode(texref, HIP_TR_FILTER_MODE_POINT)); + HIP_CHECK(hipTexRefSetFlags(texref, CU_TRSF_READ_AS_INTEGER)); + HIP_CHECK(hipTexRefSetFormat(texref, format, 1)); + HIP_CHECK(hipTexRefSetArray(texref, *array, CU_TRSA_OVERRIDE_FORMAT)); +#else + HIP_CHECK(hipModuleGetTexRef(&texref, Module, texRefName)); + HIP_CHECK(hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap)); + HIP_CHECK(hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap)); + HIP_CHECK(hipTexRefSetFilterMode(texref, hipFilterModePoint)); + HIP_CHECK(hipTexRefSetFlags(texref, HIP_TRSF_READ_AS_INTEGER)); + HIP_CHECK(hipTexRefSetFormat(texref, format, 1)); + HIP_CHECK(hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT)); +#endif +} + +template bool validateOutput(unsigned int width, + unsigned int height, + T* hData, + T* hOutputData) { + for (unsigned int i = 0; i < height; i++) { + for (unsigned int j = 0; j < width; j++) { + REQUIRE(hData[i * width + j] == hOutputData[i * width + j]); + } + } + return true; +} +/** + * Validates texture type data functionality for hipModuleGetTexRef + * + */ +template bool testTexType(hipArray_Format format, + const char* texRefName, + const char* kerFuncName) { + bool TestPassed = true; + unsigned int width = WIDTH; + unsigned int height = HEIGHT; + unsigned int size = width * height * sizeof(T); + T* hData = reinterpret_cast(malloc(size)); + if (NULL == hData) { + INFO("Info:Failed to allocate using malloc in testTexType.\n"); + return false; + } + CTX_CREATE() + fillTestBuffer(width, height, hData); + // Load Kernel File and create hipArray + hipModule_t Module; + HIP_CHECK(hipModuleLoad(&Module, CODEOBJ_FILE)); + HIP_ARRAY array; + allocInitArray(width, height, format, &array); +#if HT_NVIDIA + // Copy from hData to array using hipMemcpyParam2D + copyBuffer2Array(width, height, hData, &array); + // Get tex reference from the loaded kernel file + // Assign array to the tex reference + assignArray2TexRef(format, texRefName, Module, &array); +#else + // Copy from hData to array using hipMemcpyParam2D + copyBuffer2Array(width, height, hData, array); + // Get tex reference from the loaded kernel file + // Assign array to the tex reference + assignArray2TexRef(format, texRefName, Module, array); +#endif + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, kerFuncName)); + + T* dData = NULL; + HIP_CHECK(hipMalloc(reinterpret_cast(&dData), size)); + + struct { + void* _Ad; + unsigned int _Bd; + unsigned int _Cd; + } args; + args._Ad = reinterpret_cast(dData); + args._Bd = width; + args._Cd = height; + + size_t sizeTemp = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, + &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &sizeTemp, + HIP_LAUNCH_PARAM_END}; + + int temp1 = width / GRIDDIMX; + int temp2 = height / GRIDDIMY; + HIP_CHECK( + hipModuleLaunchKernel(Function, GRIDDIMX, GRIDDIMY, GRIDDIMZ, + temp1, temp2, BLOCKDIMZ, 0, 0, + NULL, reinterpret_cast(&config))); + HIP_CHECK(hipDeviceSynchronize()); + T* hOutputData = reinterpret_cast(malloc(size)); + if (NULL == hOutputData) { + printf("Failed to allocate using malloc in testTexType.\n"); + TestPassed = false; + } else { + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + TestPassed = validateOutput(width, height, hData, hOutputData); + } + free(hOutputData); + HIP_CHECK(hipFree(dData)); + ARRAY_DESTROY(array) + HIP_CHECK(hipModuleUnload(Module)); + free(hData); + CTX_DESTROY() + return TestPassed; +} + +/** + * Validates texture functionality with multiple streams for hipModuleGetTexRef + * + */ +template bool testTexMultStream(const std::vector& buffer, + hipArray_Format format, + const char* texRefName, + const char* kerFuncName, + unsigned int numOfStreams) { + bool TestPassed = true; + unsigned int width = WIDTH; + unsigned int height = HEIGHT; + unsigned int size = width * height * sizeof(T); + T* hData = reinterpret_cast(malloc(size)); + if (NULL == hData) { + printf("Failed to allocate using malloc in testTexMultStream.\n"); + return false; + } + CTX_CREATE() + fillTestBuffer(width, height, hData); + + // Load Kernel File and create hipArray + hipModule_t Module; + HIP_CHECK(hipModuleLoadData(&Module, &buffer[0])); + HIP_ARRAY array; + allocInitArray(width, height, format, &array); +#if HT_NVIDIA + // Copy from hData to array using hipMemcpyParam2D + copyBuffer2Array(width, height, hData, &array); + // Get tex reference from the loaded kernel file + // Assign array to the tex reference + assignArray2TexRef(format, texRefName, Module, &array); +#else + // Copy from hData to array using hipMemcpyParam2D + copyBuffer2Array(width, height, hData, array); + // Get tex reference from the loaded kernel file + // Assign array to the tex reference + assignArray2TexRef(format, texRefName, Module, array); +#endif + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, kerFuncName)); + + // Create Multiple Strings + hipStream_t streams[MAX_STREAMS]={0}; + T* dData[MAX_STREAMS] = {NULL}; + T* hOutputData[MAX_STREAMS] = {NULL}; + if (numOfStreams > MAX_STREAMS) { + numOfStreams = MAX_STREAMS; + } + unsigned int totalStreamsCreated = 0; + for (unsigned int stream_num = 0; stream_num < numOfStreams; stream_num++) { + hOutputData[stream_num] = reinterpret_cast(malloc(size)); + REQUIRE(NULL != hOutputData[stream_num]); + HIP_CHECK(hipStreamCreate(&streams[stream_num])); + HIP_CHECK(hipMalloc(reinterpret_cast(&dData[stream_num]), size)); + memset(hOutputData[stream_num], 0, size); + struct { + void* _Ad; + unsigned int _Bd; + unsigned int _Cd; + } args; + args._Ad = reinterpret_cast(dData[stream_num]); + args._Bd = width; + args._Cd = height; + + size_t sizeTemp = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, + &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &sizeTemp, + HIP_LAUNCH_PARAM_END}; + + int temp1 = width / GRIDDIMX; + int temp2 = height / GRIDDIMY; + HIP_CHECK( + hipModuleLaunchKernel(Function, GRIDDIMX, GRIDDIMY, GRIDDIMZ, + temp1, temp2, BLOCKDIMZ, 0, streams[stream_num], + NULL, reinterpret_cast(&config))); + totalStreamsCreated++; + } + // Check the kernel results separately + for (unsigned int stream_num = 0; stream_num < totalStreamsCreated; + stream_num++) { + HIP_CHECK(hipStreamSynchronize(streams[stream_num])); + HIP_CHECK(hipMemcpy(hOutputData[stream_num], dData[stream_num], size, + hipMemcpyDeviceToHost)); + TestPassed &= validateOutput(width, height, hData, + hOutputData[stream_num]); + } + for (unsigned int i = 0; i < totalStreamsCreated; i++) { + HIP_CHECK(hipFree(dData[i])); + HIP_CHECK(hipStreamDestroy(streams[i])); + free(hOutputData[i]); + } + ARRAY_DESTROY(array) + HIP_CHECK(hipModuleUnload(Module)); + free(hData); + CTX_DESTROY() + return TestPassed; +} + +/** + * Internal Thread Functions + * + */ +void launchSingleStreamMultGPU(int gpu, const std::vector& buffer) { + bool TestPassed = true; + HIP_CHECK(hipSetDevice(gpu)); + TestPassed = testTexMultStream(buffer, + HIP_AD_FORMAT_FLOAT, + "ftex", + "tex2dKernelFloat", 1); + g_thTestPassed &= static_cast(TestPassed); +} + +void launchMultStreamMultGPU(int gpu, const std::vector& buffer) { + bool TestPassed = true; + HIP_CHECK(hipSetDevice(gpu)); + TestPassed = testTexMultStream(buffer, + HIP_AD_FORMAT_FLOAT, + "ftex", + "tex2dKernelFloat", 3); + g_thTestPassed &= static_cast(TestPassed); +} +/** + * Validates texture functionality with Multiple Streams on multuple GPU + * for hipModuleGetTexRef + * + */ +bool testTexMultStreamMultGPU(int numOfGPUs, + const std::vector& buffer) { + bool TestPassed = true; + std::thread T[numOfGPUs]; + + for (int gpu = 0; gpu < numOfGPUs; gpu++) { + T[gpu] = std::thread(launchMultStreamMultGPU, gpu, buffer); + } + for (int gpu = 0; gpu < numOfGPUs; gpu++) { + T[gpu].join(); + } + + REQUIRE(TestPassed == g_thTestPassed); + return TestPassed; +} +/** + * Validates texture functionality with Single Stream on multuple GPU + * for hipModuleGetTexRef + * + */ +bool testTexSingleStreamMultGPU(int numOfGPUs, + const std::vector& buffer) { + bool TestPassed = true; + std::thread T[numOfGPUs]; + + for (int gpu = 0; gpu < numOfGPUs; gpu++) { + T[gpu] = std::thread(launchSingleStreamMultGPU, gpu, buffer); + } + for (int gpu = 0; gpu < numOfGPUs; gpu++) { + T[gpu].join(); + } + + REQUIRE(TestPassed == g_thTestPassed); + return TestPassed; +} + +TEST_CASE("Unit_hipModuleTexture2dDrv") { + bool TestPassed = true; + SECTION("testTexType is float") { + REQUIRE(TestPassed == testTexType(HIP_AD_FORMAT_FLOAT, + "ftex", "tex2dKernelFloat")); + } + SECTION("testTexType is int") { + REQUIRE(TestPassed == testTexType(HIP_AD_FORMAT_SIGNED_INT32, + "itex", "tex2dKernelInt")); + } + SECTION("testTexType is short") { + REQUIRE(TestPassed == testTexType(HIP_AD_FORMAT_SIGNED_INT16, + "stex", "tex2dKernelInt16")); + } + SECTION("testTexType is char") { + REQUIRE(TestPassed == testTexType(HIP_AD_FORMAT_SIGNED_INT8, + "ctex", "tex2dKernelInt8")); + } + SECTION("testTexMultStream is float") { + auto buffer = load_file(); + REQUIRE(TestPassed ==testTexMultStream(buffer, + HIP_AD_FORMAT_FLOAT, "ftex", "tex2dKernelFloat", MAX_STREAMS)); + } + #if HT_AMD + SECTION("testTexSingleStreamMultGPU") { + int gpu_cnt = 0; + auto buffer = load_file(); + HIP_CHECK(hipGetDeviceCount(&gpu_cnt)); + REQUIRE(TestPassed == testTexSingleStreamMultGPU(gpu_cnt, buffer)); + } + SECTION("testTexMultStreamMultGPU") { + int gpu_cnt = 0; + auto buffer = load_file(); + HIP_CHECK(hipGetDeviceCount(&gpu_cnt)); + REQUIRE(TestPassed == testTexMultStreamMultGPU(gpu_cnt, buffer)); + } + #endif + SECTION("testTexRefEqNullPtr") { + REQUIRE(TestPassed == testTexRefEqNullPtr()); + } + SECTION("testNameEqNullPtr") { + REQUIRE(TestPassed == testNameEqNullPtr()); + } + SECTION("testInvalidTexName") { + REQUIRE(TestPassed == testInvalidTexName()); + } + SECTION("testEmptyTexName") { + REQUIRE(TestPassed == testEmptyTexName()); + } + SECTION("testWrongTexRef") { + REQUIRE(TestPassed == testWrongTexRef()); + } + SECTION("testUnloadedMod") { + REQUIRE(TestPassed == testUnloadedMod()); + } +} diff --git a/catch/unit/module/hipOpenCLCOTest.cc b/catch/unit/module/hipOpenCLCOTest.cc new file mode 100644 index 000000000..a44e4dde8 --- /dev/null +++ b/catch/unit/module/hipOpenCLCOTest.cc @@ -0,0 +1,232 @@ +/* +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. +*/ + +#ifdef __linux__ +#include +#endif +#include +#include +#include + +#define OPENCL_OBJ_FILE "opencl_add.cc" +#define HIP_CODEOBJ_FILE_DEFAULT "opencl_add.co" +#define HIP_CODEOBJ_FILE_V3 "opencl_add_v3.co" +#define COMMAND_LEN 256 +#define BUFFER_LEN 256 +/** + * Validates OpenCL Static Lds Code Object + * + */ +bool testStaticLdsCodeObj(const char* pCodeObjFile) { + hipDevice_t device; + hipModule_t Module; + hipFunction_t Function; + printf("Info:Executing %s \n", __func__); + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipModuleLoad(&Module, pCodeObjFile)); + HIP_CHECK(hipModuleGetFunction(&Function, Module, "add")); + + float *Ah, *Bh; + Ah = new float[BUFFER_LEN]; + Bh = new float[BUFFER_LEN]; + for (uint32_t i = 0; i < BUFFER_LEN; i++) { + Ah[i] = i * 1.0f; + Bh[i] = 0.0f; + } + + float *Ad, *Bd; + HIP_CHECK(hipMalloc(&Ad, sizeof(float) * BUFFER_LEN)); + HIP_CHECK(hipMalloc(&Bd, sizeof(float) * BUFFER_LEN)); + HIP_CHECK(hipMemcpy(Ad, Ah, sizeof(float) * BUFFER_LEN, + hipMemcpyHostToDevice)); + + struct { + void* _Bd; + void* _Ad; + } args; + args._Ad = static_cast(Ad); + args._Bd = static_cast(Bd); + size_t size = sizeof(args); + + void *config[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END + }; + + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, BUFFER_LEN, 1, 1, 0, 0, + NULL, reinterpret_cast(&config))); + HIP_CHECK(hipMemcpy(Bh, Bd, sizeof(float) * BUFFER_LEN, + hipMemcpyDeviceToHost)); + + bool TestPassed = true; + for (uint32_t i = 0; i < BUFFER_LEN; i++) { + if (Ah[i] != Bh[i]) { + TestPassed = false; + break; + } + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + delete[] Ah; + delete[] Bh; + return TestPassed; +} + +#ifdef __linux__ +/** + * Check if environment variable $ROCM_PATH is defined + * + */ +bool isRocmPathSet() { + FILE *fpipe; + char const *command = "echo $ROCM_PATH"; + fpipe = popen(command, "r"); + + if (fpipe == nullptr) { + printf("Unable to create command\n"); + return false; + } + char command_op[BUFFER_LEN]; + if (fgets(command_op, BUFFER_LEN, fpipe)) { + size_t len = strlen(command_op); + if (len > 1) { // This is because fgets always adds newline character + pclose(fpipe); + return true; + } + } + pclose(fpipe); + return false; +} +/** + * Gets the sramecc/xnack settings from rocm info + * + */ +int getV3TargetIdFeature(char* feature, bool rocmPathSet) { + FILE *fpipe; + char command[COMMAND_LEN] = ""; + const char *rocmpath = nullptr; + if (rocmPathSet) { + // For STG2 testing where /opt/rocm path is not present + rocmpath = "$ROCM_PATH/bin/rocminfo"; + } else { + // Check if the rocminfo tool exists + rocmpath = "/opt/rocm/bin/rocminfo"; + } + snprintf(command, COMMAND_LEN, "%s", rocmpath); + strncat(command, " | grep -m1 \"sramecc.:xnack.\"", COMMAND_LEN); + fpipe = popen(command, "r"); + + if (fpipe == nullptr) { + printf("Unable to create command file\n"); + return -1; + } + char command_op[BUFFER_LEN]; + const char* pOpt1 = nullptr; + const char *pOpt2 = nullptr; + if (fgets(command_op, BUFFER_LEN, fpipe)) { + if (strstr(command_op, "sramecc+")) { + pOpt1 = "-msram-ecc"; + } else if (strstr(command_op, "sramecc-")) { + pOpt1 = "-mno-sram-ecc"; + } else { + pclose(fpipe); + return -1; + } + if (strstr(command_op, "xnack+")) { + pOpt2 = " -mxnack"; + } else if (strstr(command_op, "xnack-")) { + pOpt2 = " -mno-xnack"; + } else { + pclose(fpipe); + return -1; + } + } else { + printf("No sramecc/xnack settings found.\n"); + pclose(fpipe); + return -1; + } + strncpy(feature, pOpt1, strlen(pOpt1)); + strncat(feature, pOpt2, strlen(pOpt2)); + pclose(fpipe); + return 0; +} +#endif + +TEST_CASE("Unit_hipOpenCLCOTest") { + int version = HIP_VERSION_MAJOR; +#ifdef __linux__ + char command[COMMAND_LEN] = ""; + char v3option[32] = ""; + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + + // Generate the command to translate the OpenCL code object to hip code object + const char *pCodeObjVer = nullptr; + const char *pCodeObjFile = nullptr; + bool rocmPathSet = isRocmPathSet(); + SECTION("Default codeobj file") { + pCodeObjVer = ""; + pCodeObjFile = HIP_CODEOBJ_FILE_DEFAULT; + } + SECTION("V3 codeobj file") { + if (version >= 4) { + pCodeObjVer = "-mcode-object-version=3"; + } + if (-1 == getV3TargetIdFeature(v3option, rocmPathSet)) { + INFO("Error getting V3 Option. Skipping Test. \n"); + } + pCodeObjFile = HIP_CODEOBJ_FILE_V3; + } + /* The command string is created using multiple concatenation instead of one go + to avoid the following cpplint error: + " Multi-line string ("...") found. This lint script doesn't do well with such strings, + and may give bogus warnings. Use C++11 raw strings or concatenation instead." + */ + if (rocmPathSet) { + // For STG2 testing where /opt/rocm path is not present + snprintf(command, COMMAND_LEN, + "$ROCM_PATH/llvm/bin/clang -target amdgcn-amd-amdhsa -x cl "); + } else { + snprintf(command, COMMAND_LEN, + "/opt/rocm/llvm/bin/clang -target amdgcn-amd-amdhsa -x cl "); + } + char command_temp[COMMAND_LEN] = ""; + snprintf(command_temp, COMMAND_LEN, + "-include `find /opt/rocm* -name opencl-c.h` %s %s -mcpu=%s -o %s %s", + pCodeObjVer, v3option, props.gcnArchName, pCodeObjFile, OPENCL_OBJ_FILE); + + strncat(command, command_temp, COMMAND_LEN); + printf("command executed = %s\n", command); + + system((const char*)command); + // Check if the code object file is created + snprintf(command, COMMAND_LEN, "./%s", + pCodeObjFile); + + if (access(command, F_OK) == -1) { + INFO("Code Object File not found \n"); + } + bool TestPassed = testStaticLdsCodeObj(pCodeObjFile); + REQUIRE(TestPassed); +#else + INFO("This test is skipped due to non linux environment.\n"); +#endif +} + diff --git a/catch/unit/module/kernel_composite_test.cc b/catch/unit/module/kernel_composite_test.cc new file mode 100644 index 000000000..a66b26175 --- /dev/null +++ b/catch/unit/module/kernel_composite_test.cc @@ -0,0 +1,41 @@ +/* +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. +*/ + +#include "hip/hip_runtime.h" +#define GLOBAL_BUF_SIZE 2048 + +__device__ float deviceGlobalFloat; +__device__ int deviceGlobalInt1; +__device__ int deviceGlobalInt2; +__device__ short deviceGlobalShort; +__device__ char deviceGlobalChar; + +__device__ int getSquareOfGlobalFloat() { + return static_cast(deviceGlobalFloat*deviceGlobalFloat); +} + +extern "C" __global__ void testWeightedCopy(int* a, int* b) { + int tx = threadIdx.x; + b[tx] = deviceGlobalInt1*a[tx] + deviceGlobalInt2 + + static_cast(deviceGlobalShort) + static_cast(deviceGlobalChar) + + getSquareOfGlobalFloat(); +} diff --git a/catch/unit/module/opencl_add.cc b/catch/unit/module/opencl_add.cc new file mode 100644 index 000000000..82ff17ecb --- /dev/null +++ b/catch/unit/module/opencl_add.cc @@ -0,0 +1,37 @@ +/* +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. +*/ + +kernel void add(global float* output, global float* input) { + __local float lds[100]; + int id = get_global_id(0); + + if (id == 0) { + for (int i = 0; i < 100; i++) { + lds[i] = input[i]; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + if (id < 100) { + output[id] = lds[id]; + } else { + output[id] = input[id]; + } +} diff --git a/catch/unit/module/tex2d_kernel.cc b/catch/unit/module/tex2d_kernel.cc new file mode 100644 index 000000000..e626c80a9 --- /dev/null +++ b/catch/unit/module/tex2d_kernel.cc @@ -0,0 +1,73 @@ +/* +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. +*/ +#include "hip/hip_runtime.h" + +texture ftex; +texture itex; +texture stex; +texture ctex; + +__device__ float deviceGlobalFloat; + +extern "C" __global__ void tex2dKernelFloat(float* outputData, + int width, int height) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if ((x < width) && (y < width)) { + outputData[y * width + x] = tex2D(ftex, x, y); + } +#endif +} + +extern "C" __global__ void tex2dKernelInt(int* outputData, + int width, int height) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if ((x < width) && (y < width)) { + outputData[y * width + x] = tex2D(itex, x, y); + } +#endif +} + +extern "C" __global__ void tex2dKernelInt16(short* outputData, + int width, int height) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if ((x < width) && (y < width)) { + outputData[y * width + x] = tex2D(stex, x, y); + } +#endif +} + +extern "C" __global__ void tex2dKernelInt8(char* outputData, + int width, int height) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + if ((x < width) && (y < width)) { + outputData[y * width + x] = tex2D(ctex, x, y); + } +#endif +} diff --git a/catch/unit/module/vcpy_kernel.cc b/catch/unit/module/vcpy_kernel.cc new file mode 100644 index 000000000..299d212fa --- /dev/null +++ b/catch/unit/module/vcpy_kernel.cc @@ -0,0 +1,25 @@ +/* +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. +*/ + +#include "hip/hip_runtime.h" + +extern "C" __global__ void hello_world(float* a, float* b) { + int tx = threadIdx.x; + b[tx] = a[tx]; +}