diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 7d6fd3f731..5d4a319187 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -30,16 +30,17 @@ THE SOFTWARE. */ /** - * @defgroup GraphTest Graph Management + * @defgroup ContextTest Context Management * @{ - * This section describes tests for the graph management types & functions of HIP runtime API. + * This section describes tests for the context management functions of HIP runtime API. + * @warning All Context Management APIs are **deprecated** and shall not be implemented. * @} */ /** - * @defgroup VectorTypeTest Vector types + * @defgroup DeviceLanguageTest Device Language * @{ - * This section describes tests for the Vector type functions and operators. + * This section describes tests for the Device Language API. * @} */ @@ -50,28 +51,6 @@ THE SOFTWARE. * @} */ -/** - * @defgroup EventTest Event Management - * @{ - * This section describes tests for the event management functions of HIP runtime API. - * @} - */ - -/** - * @defgroup ErrorTest Error Handling - * @{ - * This section describes tests for the error handling functions of HIP runtime API. - * @} - */ - -/** - * @defgroup PeerToPeerTest PeerToPeer Device Memory Access - * @{ - * This section describes tests for the PeerToPeer device memory access functions of HIP runtime API. - * @warning PeerToPeer support is experimental. - * @} - */ - /** * @defgroup DriverTest Initialization and Version * @{ @@ -79,56 +58,6 @@ THE SOFTWARE. * @} */ -/** - * @defgroup DeviceLanguageTest Device Language - * @{ - * This section describes tests for the Device Language API. - * @} - */ - -/** - * @defgroup ExecutionTest Execution Control - * @{ - * This section describes tests for the execution control functions of HIP runtime API. - * @} - */ - -/** - * @defgroup TextureTest Texture Management - * @{ - * This section describes tests for the texture management functions of HIP runtime API. - * @} - */ - - /** - * @defgroup ShflTest warp shuffle function Management - * @{ - * This section describes the warp shuffle types & functions of HIP runtime API. - * @} - */ - -/** - * @defgroup EventTest Event Management - * @{ - * This section describes the event management types & functions of HIP runtime API. - * @} - */ - -/** - * @defgroup ContextTest Context Management - * @{ - * This section describes tests for the context management functions of HIP runtime API. - * @warning All Context Management APIs are **deprecated** and shall not be implemented. - * @} - */ - -/** - * @defgroup StreamTest Stream Management - * @{ - * This section describes the stream management types & functions of HIP runtime API. - * @} - */ - /** * @defgroup DynamicLoadingTest Kernel Loading Management * @{ @@ -137,9 +66,30 @@ THE SOFTWARE. */ /** - * @defgroup MemoryTest memory Management APIs + * @defgroup ErrorTest Error Handling * @{ - * This section describes the memory management types & functions of HIP runtime API. + * This section describes tests for the error handling functions of HIP runtime API. + * @} + */ + +/** + * @defgroup EventTest Event Management + * @{ + * This section describes tests for the event management functions of HIP runtime API. + * @} + */ + +/** + * @defgroup ExecutionTest Execution Control + * @{ + * This section describes tests for the execution control functions of HIP runtime API. + * @} + */ + +/** + * @defgroup GraphTest Graph Management + * @{ + * This section describes tests for the graph management types & functions of HIP runtime API. * @} */ @@ -151,14 +101,44 @@ THE SOFTWARE. */ /** - * @defgroup DeviceLanguageTest Device Language + * @defgroup MemoryTest memory Management APIs * @{ - * This section describes tests for the Device Language API. + * This section describes the memory management types & functions of HIP runtime API. * @} */ /** - * @defgroup DeviceLanguageTest Device Language + * @defgroup PeerToPeerTest PeerToPeer Device Memory Access * @{ - * This section describes tests for the Device Language API. + * This section describes tests for the PeerToPeer device memory access functions of HIP runtime API. + * @warning PeerToPeer support is experimental. + * @} + */ + + /** + * @defgroup ShflTest warp shuffle function Management + * @{ + * This section describes the warp shuffle types & functions of HIP runtime API. + * @} + */ + +/** + * @defgroup StreamTest Stream Management + * @{ + * This section describes the stream management types & functions of HIP runtime API. + * @} + */ + +/** + * @defgroup TextureTest Texture Management + * @{ + * This section describes tests for the texture management functions of HIP runtime API. + * @} + */ + +/** + * @defgroup VectorTypeTest Vector types + * @{ + * This section describes tests for the Vector type functions and operators. + * @} */ diff --git a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt index 77a48363ed..d8a07b8638 100644 --- a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt @@ -23,6 +23,13 @@ set(TEST_SRC hipMemFaultStackAllocation.cc hipLaunchBounds.cc hipShflTests.cc + hipDynamicShared.cc + hipDynamicShared2.cc + hipEmptyKernel.cc + hipGridLaunch.cc + hipLanguageExtensions.cc + hipLaunchParm.cc + hipLaunchParmFunctor.cc hipTestConstant.cc hipTestGlobalVariable.cc hipTestMemKernel.cc @@ -47,6 +54,7 @@ endif() if(HIP_PLATFORM MATCHES "amd") set(AMD_SRC hipShflUpDownTest.cc + hipExtLaunchKernelGGL.cc ) set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) endif() diff --git a/projects/hip-tests/catch/unit/kernel/hipDynamicShared.cc b/projects/hip-tests/catch/unit/kernel/hipDynamicShared.cc new file mode 100644 index 0000000000..168d979d14 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipDynamicShared.cc @@ -0,0 +1,161 @@ +/* +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 +unsigned threadsPerBlock = 256; + +template +__device__ void sum(T* sdata, unsigned groupElements, unsigned tid) { + T tmp; + if (groupElements < batch) + return; + // sdata[tid] += sdata[tid - batch/2] does not work when block size is + // greater than wave size because one wave may complete before another + // wave. + if (tid >= batch/2 && tid < groupElements) + tmp = sdata[tid - batch/2]; + __syncthreads(); + if (tid >= batch/2 && tid < groupElements) + sdata[tid] += tmp; + __syncthreads(); +} + +template +__global__ void testExternSharedKernel(const T* A_d, const T* B_d, T* C_d, + size_t numElements, size_t groupElements) { + // declare dynamic shared memory + extern __shared__ double sdata0[]; + T* sdata = reinterpret_cast(sdata0); + + size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); + size_t tid = threadIdx.x; + + // initialize dynamic shared memory + if (tid < groupElements) { + sdata[tid] = static_cast(tid); + } + __syncthreads(); + + // prefix sum inside dynamic shared memory + sum<512>(sdata, groupElements, tid); + sum<256>(sdata, groupElements, tid); + sum<128>(sdata, groupElements, tid); + sum<64>(sdata, groupElements, tid); + sum<32>(sdata, groupElements, tid); + sum<16>(sdata, groupElements, tid); + sum<8>(sdata, groupElements, tid); + sum<4>(sdata, groupElements, tid); + sum<2>(sdata, groupElements, tid); + C_d[gid] = A_d[gid] + B_d[gid] + sdata[tid % groupElements]; +} + +template +void testExternShared(size_t N, unsigned groupElements) { + size_t Nbytes = N * sizeof(T); + + T *A_d, *B_d, *C_d; + T *A_h, *B_h, *C_h; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = N/threadsPerBlock; + assert(N == blocks * threadsPerBlock); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + // calculate the amount of dynamic shared memory required + size_t groupMemBytes = groupElements * sizeof(double); + + // launch kernel with dynamic shared memory + hipLaunchKernelGGL(HIP_KERNEL_NAME(testExternSharedKernel), dim3(blocks), + dim3(threadsPerBlock), groupMemBytes, 0, A_d, B_d, C_d, + N, groupElements); + + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + // verify + for (size_t i = 0; i < N; ++i) { + size_t tid = (i % min(threadsPerBlock, groupElements)); + T sumFromSharedMemory = static_cast(tid * (tid + 1) / 2); + T expected = A_h[i] + B_h[i] + sumFromSharedMemory; + REQUIRE(C_h[i] == expected); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - launch kernel with dynamic shared memory for float and double + * datatypes and verify the results. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipDynamicShared.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipDynamicShared") { + SECTION("test case with float for least size") { + testExternShared(1024, 4); + testExternShared(1024, 8); + testExternShared(1024, 16); + testExternShared(1024, 32); + testExternShared(1024, 64); + } + + SECTION("test case with float for max size") { + testExternShared(65536, 4); + testExternShared(65536, 8); + testExternShared(65536, 16); + testExternShared(65536, 32); + testExternShared(65536, 64); + } + + SECTION("test case with double for least size") { + testExternShared(1024, 4); + testExternShared(1024, 8); + testExternShared(1024, 16); + testExternShared(1024, 32); + testExternShared(1024, 64); + } + + SECTION("test case with double for max size") { + testExternShared(65536, 4); + testExternShared(65536, 8); + testExternShared(65536, 16); + testExternShared(65536, 32); + testExternShared(65536, 64); + } +} diff --git a/projects/hip-tests/catch/unit/kernel/hipDynamicShared2.cc b/projects/hip-tests/catch/unit/kernel/hipDynamicShared2.cc new file mode 100644 index 0000000000..47a94c1357 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipDynamicShared2.cc @@ -0,0 +1,89 @@ +/* +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 + +#define LEN (16 * 1024) +#define SIZE (LEN * sizeof(float)) + +__global__ void vectorAdd(float* Ad, float* Bd) { + extern __shared__ float sBd[]; + int tx = threadIdx.x; + for (int i = 0; i < LEN / 64; i++) { + sBd[tx + i * 64] = Ad[tx + i * 64] + 1.0f; + Bd[tx + i * 64] = sBd[tx + i * 64]; + } +} + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Assign max dynamic shared memory to kernel function and + * verify the results. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipDynamicShared2.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipDynamicShared2") { + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + for (int i = 0; i < LEN; i++) { + A[i] = 1.0f; + B[i] = 1.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)); + + hipError_t ret = hipFuncSetAttribute( + reinterpret_cast(&vectorAdd), + hipFuncAttributeMaxDynamicSharedMemorySize, SIZE); + + REQUIRE(ret == hipSuccess); + hipLaunchKernelGGL(vectorAdd, dim3(1, 1, 1), dim3(64, 1, 1), SIZE, 0, Ad, Bd); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + for (int i = 0; i < LEN; i++) { + assert(B[i] > 1.0f && B[i] < 3.0f); + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + + delete[] A; + delete[] B; +} diff --git a/projects/hip-tests/catch/unit/kernel/hipEmptyKernel.cc b/projects/hip-tests/catch/unit/kernel/hipEmptyKernel.cc new file mode 100644 index 0000000000..484a77e0a6 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipEmptyKernel.cc @@ -0,0 +1,52 @@ +/* +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 + +__global__ void Empty(int param) {} + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - pass empty Kernel function. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipEmptyKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipEmptyKernel") { + hipLaunchKernelGGL(HIP_KERNEL_NAME(Empty), dim3(1), dim3(1), 0, 0, 0); + HIP_CHECK(hipDeviceSynchronize()); +} diff --git a/projects/hip-tests/catch/unit/kernel/hipExtLaunchKernelGGL.cc b/projects/hip-tests/catch/unit/kernel/hipExtLaunchKernelGGL.cc new file mode 100644 index 0000000000..c23ebb7f81 --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipExtLaunchKernelGGL.cc @@ -0,0 +1,133 @@ +/* +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. +*/ +// Test the Grid_Launch syntax. + +#include +#include +#include +#include +#include "hip/hip_ext.h" + +static unsigned threadsPerBlock = 256; +static unsigned blocksPerCU = 6; + +struct _t { + double _a, _b, _c, _d, _e, _f, _g, _h, _i, _j; +}; + +typedef struct _t _T; + +__global__ void sKernel(_T s, double *a) { + *a = s._a + s._b + s._c + s._d + s._e + s._f + s._g + s._h + s._i + s._j; +} + +__global__ void mKernel(char f, int16_t a, int b, double c, + int16_t d, int e, double* res) { + *res = a + b + c + d + e + f; +} + +void testMixData() { + double m = 0; + double *d_m; + HIP_CHECK(hipMalloc(&d_m, sizeof(double))); + int a = 1, e = 10; + int16_t b = 2, d = 4; + double c = 3.0; + char ff = 10; + hipExtLaunchKernelGGL(mKernel, 1, 1, 0, 0, nullptr, nullptr, 0, ff, + b, a, c, d, e, d_m); + HIP_CHECK(hipMemcpy(&m, d_m, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(m == 30.0); + HIP_CHECK(hipFree(d_m)); +} + +void testStruct() { + double m = 0; + double *d_m; + HIP_CHECK(hipMalloc(&d_m, sizeof(double))); + _T s{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + hipExtLaunchKernelGGL(sKernel, 1, 1, 0, 0, nullptr, nullptr, 0, s, d_m); + HIP_CHECK(hipMemcpy(&m, d_m, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(m == 55.0); + HIP_CHECK(hipFree(d_m)); +} + +void test(size_t N) { + size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipExtLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), + dim3(threadsPerBlock), 0, 0, nullptr, nullptr, 0, + static_cast(A_d), + static_cast(B_d), C_d, N); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); +} + +/** +* @addtogroup hipExtLaunchKernelGGL hipExtLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, + hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags, + Args... args)` - +* Launches kernel with dimention parameters and shared memory on stream with templated kernel and arguments +*/ + +/** + * Test Description + * ------------------------ + * - Test case to verify sample array with hipExtLaunchKernelGGL() + * and verify the results. + * - Test case to verify struct data with hipExtLaunchKernelGGL() + * and verify the results. + * - Test case to verify mix datatypes with hipExtLaunchKernelGGL() + * and verify the results. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipExtLaunchKernelGGL.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipExtLaunchKernelGGL") { + SECTION("test run") { + size_t N = 4 * 1024 * 1024; + test(N); + } + SECTION("testStruct run") { + testStruct(); + } + SECTION("testMixData run") { + testMixData(); + } +} diff --git a/projects/hip-tests/catch/unit/kernel/hipGridLaunch.cc b/projects/hip-tests/catch/unit/kernel/hipGridLaunch.cc new file mode 100644 index 0000000000..7716b0781a --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipGridLaunch.cc @@ -0,0 +1,117 @@ +/* +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. +*/ +// Test the Grid_Launch syntax. + +#include +#include +#include +#include + +static unsigned threadsPerBlock = 256; +static unsigned blocksPerCU = 6; + +// __device__ maps to __attribute__((hc)) +__device__ int foo(int i) { return i + 1; } + + +template +__global__ void vectorADD2(T* A_d, T* B_d, T* C_d, size_t N) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < N; i += stride) { + double foo = __hiloint2double(A_d[i], B_d[i]); + C_d[i] = __double2loint(foo) + __double2hiint(foo); + } +} + +int test_gl2(size_t N) { + size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + // Full vadd in one large chunk, to get things started: + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(vectorADD2, dim3(blocks), dim3(threadsPerBlock), + 0, 0, A_d, B_d, C_d, N); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + // verify + HipTest::checkVectorADD(A_h, B_h, C_h, N); + return 0; +} + +#if __HIP__ +int test_triple_chevron(size_t N) { + size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + // Full vadd in one large chunk, to get things started: + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + vectorADD2<<>>(A_d, B_d, C_d, N); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + // verify + HipTest::checkVectorADD(A_h, B_h, C_h, N); + return 0; +} +#endif + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test case to verify the Grid_Launch syntax. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipGridLaunch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipGridLaunch") { + size_t N = 4 * 1024 * 1024; + SECTION("Test test_gl2") { + test_gl2(N); + } + +#if __HIP__ + SECTION("Test triple_chevron") { + test_triple_chevron(N); + } +#endif +} diff --git a/projects/hip-tests/catch/unit/kernel/hipLanguageExtensions.cc b/projects/hip-tests/catch/unit/kernel/hipLanguageExtensions.cc new file mode 100644 index 0000000000..5e601d35cf --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipLanguageExtensions.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 + +// Simple tests for variable type qualifiers: +__device__ int deviceVar; + +// TODO-HCC __constant__ not working yet. +__constant__ int constantVar1; + +__constant__ __device__ int constantVar2; + +// Test HOST space: +__host__ void foo() { printf("foo!\n"); } + +__device__ __noinline__ int sum1_noinline(int a) { return a + 1; } +__device__ __forceinline__ int sum1_forceinline(int a) { return a + 1; } + + +__device__ __host__ float PlusOne(float x) { return x + 1.0; } + +__global__ void MyKernel(const float* a, const float* b, float* c, + unsigned N) { + unsigned gid = threadIdx.x; + if (gid < N) { + c[gid] = a[gid] + PlusOne(b[gid]); + } +} + +void callMyKernel() { + float *a, *b, *c; + const unsigned blockSize = 256; + unsigned N = blockSize; + hipLaunchKernelGGL(MyKernel, dim3(N / blockSize), dim3(blockSize), + 0, 0, a, b, c, N); +} + +template +__global__ void vectorADD(T __restrict__* A_d, T* B_d, T* C_d, size_t N) { +#ifdef NOT_YET + int a = __shfl_up(x, 1); +#endif + float x = 1.0; + float z = sin(x); +#ifdef NOT_YET + float fastZ = __sin(x); +#endif + __syncthreads(); + + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < N; i += stride) { + C_d[i] = A_d[i] + B_d[i]; + } +} + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Collection of code to make sure that various features + * in the hip kernel language compile. + + * Test source + * ------------------------ + * - catch/unit/kernel/hipLanguageExtensions.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipLanguageExtensions") { + REQUIRE(true); +} diff --git a/projects/hip-tests/catch/unit/kernel/hipLaunchParm.cc b/projects/hip-tests/catch/unit/kernel/hipLaunchParm.cc new file mode 100644 index 0000000000..f033829f6e --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipLaunchParm.cc @@ -0,0 +1,1009 @@ +/* +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 + +// Memory alignment is broken +// Update: with latest changes the aligment is working fine, hence enabled +#define ENABLE_ALIGNMENT_TEST_SMALL_BAR 1 + +// Packed member atribute broken +#define ENABLE_PACKED_TEST 0 + +// Update: with latest changes struct class object +// from device is working fine, hence enabled +#define ENABLE_CLASS_OBJ_ACCESS 1 + +// accessing dynamic/heap memory from device is broken +#define ENABLE_HEAP_MEMORY_ACCESS 0 + +// Update: with latest changes it's working hence enabled +#define ENABLE_USER_STL 1 + +// Update: with latest changes it's working hence enabled +#define ENABLE_OUT_OF_ORDER_INITIALIZATION 1 + +// Direct initialization of struct broken, +// ip_d9 is a pointer, uint_t*, hipLaunchKernelStruct_h9 = {'c', ip_d9}; +#define ENABLE_DECLARE_INITIALIZATION_POINTER 0 + +// Bit fields are broken +#define ENABLE_BIT_FIELDS 0 + +static const int BLOCK_DIM_SIZE = 512; + +// allocate memory on device and host for result validation +static bool *result_d, *result_h; + +static hipError_t hipMallocError = hipErrorUnknown; +static hipError_t hipHostMallocError = hipErrorUnknown; +static hipError_t hipMemsetError = hipErrorUnknown; + +static void ResultValidation() { + HIP_CHECK(hipMemcpy(result_h, result_d, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + REQUIRE(result_h[k] == true); + } + return; +} + +// Segregating the reset part as it was causing a problem when i put inside +// ResultValidation() function, the memory was not reset correctly for the +// tests which were disabled. +static void ResetValidationMem() { + // reset the memory to false to reuse it. + HIP_CHECK(hipMemset(result_d, false, BLOCK_DIM_SIZE)); + HIP_CHECK(hipMemset(result_h, false, BLOCK_DIM_SIZE)); + return; +} + +// This test is to verify Struct with variables +// support, read from device. +typedef struct hipLaunchKernelStruct1 { + int li; // local int + float lf; // local float + bool result; // local bool +} hipLaunchKernelStruct_t1; + +// This test is to verify struct with padding, read from device +typedef struct hipLaunchKernelStruct2 { + char c1; + int64_t l1; + char c2; + int64_t l2; + bool result; +} hipLaunchKernelStruct_t2; + +// This test is to verify struct with padding, read from device +typedef struct hipLaunchKernelStruct3 { + char bf1; + char bf2; + int64_t l1; + char bf3; + bool result; +} hipLaunchKernelStruct_t3; + +// This test is to verify empty struct +typedef struct hipLaunchKernelStruct4 { + // empty struct, size will be verified from device side,size 1Byte +} hipLaunchKernelStruct_t4; + +// This test is to verify struct with pointer member variable. +typedef struct hipLaunchKernelStruct5 { + char c1; + char* cp; // char pointer +} hipLaunchKernelStruct_t5; + + +// This test is to verify struct with aligned(8), +// right now it's broken on hcc & hip-clang +typedef struct hipLaunchKernelStruct6 { + char c1; + int16_t si; +} __attribute__((aligned(8))) hipLaunchKernelStruct_t6; + +// This test is to verify struct with aligned(16), +// right now it's brokenon hcc & hip-clang +typedef struct hipLaunchKernelStruct7 { + char c1; + int16_t si; +} __attribute__((aligned(16))) hipLaunchKernelStruct_t7; + +// This test is to verify struct with packed & aligned, +// size should be 4Bytes right now it's broken on hcc & hip-clang +typedef struct hipLaunchKernelStruct8 { + char c1; + int16_t si; + bool b; +}__attribute__((packed, aligned(4))) hipLaunchKernelStruct_t8; + +// This test is to verify struct with packed, no alignment as Sam suggested +// size should be 4Bytes, right now it's broken on hcc & hip-clang +typedef struct hipLaunchKernelStruct8A { + char c1; + int16_t si; + bool b; +}__attribute__((packed)) hipLaunchKernelStruct_t8A; + +// This test is to verify struct with alignment, no packing as Sam suggested +// size should be 8Bytes as no packing, right now it's broken on hcc & hip-clang +typedef struct hipLaunchKernelStruct8B { + char c1; + int16_t si; + bool b; +}__attribute__((aligned(8))) hipLaunchKernelStruct_t8B; + +// This test is to verify const struct object +typedef struct hipLaunchKernelStruct9 { + char c1; + uint32_t* ip; // uint pointer +} hipLaunchKernelStruct_t9; + +// This test is to verify struct with stdint types, uintN_t +typedef struct hipLaunchKernelStruct10 { + uint64_t u64; + uint32_t u32; + uint8_t u8; +} hipLaunchKernelStruct_t10; + +// This test is to verify struct with volatile member +typedef struct hipLaunchKernelStruct11 { + int i1; + volatile unsigned int vint; +} hipLaunchKernelStruct_t11; + +// This test is to verify struct with simple class object +class base { + public: + int i = 0; + base() {} +}; +typedef struct hipLaunchKernelStruct12 { + base b; + char c1; +} hipLaunchKernelStruct_t12; + +// This test is to verify struct with __device__ func() attribute +typedef struct hipLaunchKernelStruct13 { + int i1; + __device__ int getvalue() { return i1; } +} hipLaunchKernelStruct_t13; + +// This test is to verify struct with array variable, +// write to from device +typedef struct hipLaunchKernelStruct14 { + int readint; + int writeint[BLOCK_DIM_SIZE]; // will write to this from device +} hipLaunchKernelStruct_t14; + +// This test is to verify struct with dynamic memory, new int +// the heap memory will be accessed from device +typedef struct hipLaunchKernelStruct15 { + char c1; + int* heapmem; // allocated using hipMalloc() +} hipLaunchKernelStruct_t15; + +// This test is to verify simple template struct +template +struct hipLaunchKernelStruct_t16 { + T t1; +}; + +// This test is to verify simple explicity template struct +template struct hipLaunchKernelStruct_t17 {}; +template<> // explicit template +struct hipLaunchKernelStruct_t17 { + int t1; +}; + +// This test is to verity write to struct memory using __device__ func() +typedef struct hipLaunchKernelStruct18 { + char c1; + __device__ void setChar(char c) { c1 = c; } + __device__ int getChar() { return c1; } +} hipLaunchKernelStruct_t18; + +// This test is to verity user defined STL, simple stack implementation +typedef struct stackNode { + int data; + stackNode* nextNode = NULL; +} stackNode_t; +typedef struct hipLaunchKernelStruct19 { + stackNode_t* stack = NULL; + unsigned int size_ = 0; + void pushMe(int value) { // not a device function, setting from host + stackNode_t* newNode; + HIP_CHECK(hipMalloc(reinterpret_cast(&newNode), + sizeof(stackNode_t))); + HIP_CHECK(hipMemset(&newNode->data, value, sizeof(stackNode_t))); + // newNode->data = value; + ++size_; + if (stack == NULL) { + stack = newNode; + return; + } + stackNode_t* currentHead = stack; + stack = newNode; + stack->nextNode = currentHead; + return; + } + __device__ void popMe() { + stackNode_t* currentHead = stack; + stack = stack->nextNode; + --size_; + // delete currentHead; // no idea why delete not working + return; + } + int stackSize() { + return size_; + } +} hipLaunchKernelStruct_t19; + +// This test is to verify out of order initalizer of struct elements +// and access in-order, from device. +typedef struct hipLaunchKernelStruct20 { + char name; + int age; + int rank; +} hipLaunchKernelStruct_t20; + +// This test is to verify bit fields operations +// the size should be 1Bytes +typedef struct hipLaunchKernelStruct21 { + int i : 3; // limiting bits to 3 + int j : 2; // limiting bits to 2 +} hipLaunchKernelStruct_t21; + +// Passing struct to a hipLaunchKernelGGL(), +// read and write into the same struct +__global__ void hipLaunchKernelStructFunc1( + hipLaunchKernelStruct_t1 hipLaunchKernelStruct_, + bool* result_d1) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d1[x] = ((hipLaunchKernelStruct_.li == 1) + && (hipLaunchKernelStruct_.lf == 1.0) + && (hipLaunchKernelStruct_.result == false)); +} + +// Passing struct to a hipLaunchKernelGGL(), checks padding, +// read and write into the same struct +__global__ void hipLaunchKernelStructFunc2( + hipLaunchKernelStruct_t2 hipLaunchKernelStruct_, + bool* result_d2) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d2[x] = ((hipLaunchKernelStruct_.c1 == 'a') + && (hipLaunchKernelStruct_.l1 == 1.0) + && (hipLaunchKernelStruct_.c2 == 'b') + && (hipLaunchKernelStruct_.l2 == 2.0) ); +} + +// Passing struct to a hipLaunchKernelGGL(), checks padding, +// read and write into the same struct +__global__ void hipLaunchKernelStructFunc3( + hipLaunchKernelStruct_t3 hipLaunchKernelStruct_, + bool* result_d3) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d3[x] = ((hipLaunchKernelStruct_.bf1 == 1) + && (hipLaunchKernelStruct_.bf2 == 1) + && (hipLaunchKernelStruct_.l1 == 1.0) + && (hipLaunchKernelStruct_.bf3 == 1) ); +} + +// Passing empty struct to a hipLaunchKernelGGL(), +// check the size of 1Byte, set result_d4 to true if condition met +__global__ void hipLaunchKernelStructFunc4( + hipLaunchKernelStruct_t4 hipLaunchKernelStruct_, + bool* result_d4) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d4[x] = (sizeof(hipLaunchKernelStruct_) == 1); +} + +// Passing struct with pointer object to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc5( + hipLaunchKernelStruct_t5 hipLaunchKernelStruct_, + bool* result_d5) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d5[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (*hipLaunchKernelStruct_.cp == 'p')); +} + +// Passing struct which is aligned to 8Byte to a hipLaunchKernelGGL(), +// set the result_d6 to true if condition met +__global__ void hipLaunchKernelStructFunc6( + hipLaunchKernelStruct_t6 hipLaunchKernelStruct_, + bool* result_d6) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + // get the address of the struct + // size_t(p)%8 will be 0 if aligned to 8Byte address space + int *p = reinterpret_cast(&hipLaunchKernelStruct_); + result_d6[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%8 ==0)); +} + +// Passing struct which is aligned to 16Byte, +// set the result_d7 to true if condition met +__global__ void hipLaunchKernelStructFunc7( + hipLaunchKernelStruct_t7 hipLaunchKernelStruct_, + bool* result_d7) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + // get the address of the struct + // size_t(p)%16 will be 0 if aligned to 16Byte address space + int *p = reinterpret_cast(&hipLaunchKernelStruct_); + result_d7[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%16 ==0) ); +} + +// Passing struct which is packed & aligned to 4Byte, +// set the result_d8 to true if condition met +__global__ void hipLaunchKernelStructFunc8( + hipLaunchKernelStruct_t8 hipLaunchKernelStruct_, + bool* result_d8) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + // get the address of the xth element, struct[x], + // size_t(p)%4 will be 0 if aligned to 4Byte address space + int *p = reinterpret_cast(&hipLaunchKernelStruct_); + result_d8[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%4 ==0) + && (sizeof(hipLaunchKernelStruct_) == 4)); +} + +// Passing struct which is packed only, as Sam suggested, should be 4Bytes +// set the result_d8A to true if condition met +__global__ void hipLaunchKernelStructFunc8A( + hipLaunchKernelStruct_t8A hipLaunchKernelStruct_, + bool* result_d8A) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + // this is packed struct + // the address will not be aglined in this case hence condition removed + // only sizeof(hipLaunchKernelStruct_) will be valided + result_d8A[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && (sizeof(hipLaunchKernelStruct_) == 4)); +} + +// Passing struct which is aligned(4) only, as Sam suggested +// , size should be 8Bytes, set the result_d8B to true if condition met +__global__ void hipLaunchKernelStructFunc8B( + hipLaunchKernelStruct_t8B hipLaunchKernelStruct_, + bool* result_d8B) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + // get the address of the xth element, struct[x], + // size_t(p)%4 will be 0 if aligned to 4Byte address space + int *p = reinterpret_cast(&hipLaunchKernelStruct_); + result_d8B[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%8 == 0) + && (sizeof(hipLaunchKernelStruct_) == 8)); +} + +// Passing struct with uint pointer object to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc9( + const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_, + bool* result_d9) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d9[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (*hipLaunchKernelStruct_.ip == 1)); +} + +// Passing struct with stdint types object, uintN_t, to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc10( + hipLaunchKernelStruct_t10 hipLaunchKernelStruct_, + bool* result_d10) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d10[x] = ((hipLaunchKernelStruct_.u64 == UINT64_MAX) + && (hipLaunchKernelStruct_.u32 == 1) + && (hipLaunchKernelStruct_.u8 == UINT8_MAX)); +} + +// Passing struct with volatile member, to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc11( + hipLaunchKernelStruct_t11 hipLaunchKernelStruct_, + bool* result_d11) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d11[x] = ((hipLaunchKernelStruct_.i1 == 1) + && (hipLaunchKernelStruct_.vint == 0)); +} + +// Passing struct with simple class obj, to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc12( + hipLaunchKernelStruct_t12 hipLaunchKernelStruct_, + bool* result_d12) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d12[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.b.i == 0)); +} + +// Passing struct with simple __device__ func(), to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc13( + hipLaunchKernelStruct_t13 hipLaunchKernelStruct_, + bool* result_d13) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d13[x] = ((hipLaunchKernelStruct_.i1 == 1) + && (hipLaunchKernelStruct_.getvalue() == 1)); +} + +// Passing struct with array variable, write to from device +__global__ void hipLaunchKernelStructFunc14( + hipLaunchKernelStruct_t14 hipLaunchKernelStruct_, + bool* result_d14) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + hipLaunchKernelStruct_.writeint[x] = 1; + // set the result to true if the condition met + result_d14[x] = ((hipLaunchKernelStruct_.readint == 1) + && (hipLaunchKernelStruct_.writeint[x] == 1)); +} + +// Passing struct with struct with dynamic memory, new int +// the heap memory will be accessed from device +__global__ void hipLaunchKernelStructFunc15( + hipLaunchKernelStruct_t15 hipLaunchKernelStruct_, + bool* result_d15) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d15[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.heapmem[x] == 1)); +} + +// Passing simple template struct +__global__ void hipLaunchKernelStructFunc16( + hipLaunchKernelStruct_t16 hipLaunchKernelStruct_, + bool* result_d16) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d16[x] = (hipLaunchKernelStruct_.t1 == 'c'); +} + +// Passing simple explicit template struct +__global__ void hipLaunchKernelStructFunc17( + hipLaunchKernelStruct_t17 hipLaunchKernelStruct_, + bool* result_d17) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // set the result to true if the condition met + result_d17[x] = (hipLaunchKernelStruct_.t1 == 1); +} + +// Passing struct and write to struct memory using __device__ func() +__global__ void hipLaunchKernelStructFunc18( + hipLaunchKernelStruct_t18 hipLaunchKernelStruct_, + bool* result_d18) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + hipLaunchKernelStruct_.setChar('c'); + // set the result to true if the condition met + result_d18[x] = (hipLaunchKernelStruct_.getChar() == 'c'); +} + +// Passing simple user defined stack implemenration, using __device__ func() +__global__ void hipLaunchKernelStructFunc19( + hipLaunchKernelStruct_t19 hipLaunchKernelStruct_) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // stack should be empty after the kernel execustion, verify on host side + hipLaunchKernelStruct_.popMe(); +} + +// Passing out of order initalized struct, access in-order +__global__ void hipLaunchKernelStructFunc20( + hipLaunchKernelStruct_t20 hipLaunchKernelStruct_, + bool* result_d20) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // accessing struct members in order + result_d20[x] = (hipLaunchKernelStruct_.name == 'A' + // strcmp(hipLaunchKernelStruct_.name, "AMD") -> strcmp is not broken + && hipLaunchKernelStruct_.age == 42 + && hipLaunchKernelStruct_.rank == 2); +} + +// Passing struct with bit fields +__global__ void hipLaunchKernelStructFunc21( + hipLaunchKernelStruct_t21 hipLaunchKernelStruct_, + bool* result_d21) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + // accessing struct members in order + result_d21[x] = (hipLaunchKernelStruct_.i == 2 + && hipLaunchKernelStruct_.j == 0 + && (sizeof(hipLaunchKernelStruct_) == 1)); +} + +__global__ void vAdd(float* a) {} + +template +__global__ void myKernel(T1 a, T2 b) {} + + +//--- +// Some wrapper macro for testing: +#define WRAP(...) __VA_ARGS__ + +#define MY_LAUNCH_MACRO(cmd, elapsed, quiet) \ + do { \ + HIP_CHECK(hipDeviceSynchronize()); \ + cmd; \ + HIP_CHECK(hipDeviceSynchronize()); \ + } while (0); + + +#define MY_LAUNCH(command, doTrace, msg) \ + { \ + if (doTrace) printf("TRACE: %s %s\n", msg, #command); \ + command; \ + } + + +#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \ + { \ + if (doTrace) printf("TRACE: %s %s\n", msg, #command); \ + (command); \ + } + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Passing struct to a hipLaunchKernelGGL(), + * read and write into the same struct + * - Test to verify by Passing Struct type, checks padding + * - Test to verify by Passing Struct type, checks padding, assigning integer to a char + * - Test to verify by Passing empty struct + * - Test to verify by Passing struct with pointer object to a hipLaunchKernelGGL() + * - Test to verify by Passing struct with aligned(8) + * - Test to verify by Passing struct with aligned(16) + * - Test to verify by Passing struct with packed aligned to 4Bytes + * - Test to verify by Passing struct with packed to 4Bytes + * - Test to verify by Passing struct with aligned(4) to 4Bytes, size is 8Bytes + * - Test to verify by Passing const struct object to a hipLaunchKernelGGL() + * - Test to verify by Passing struct with uintN_t as member variables + * - Test to verify by Passing struct with uintN_t as member variables + * - Test to verify by Passing struct with simple class object + * - Test to verify by Passing struct with simple __device__ func() + * - Test to verify by Passing struct with array variable, write to from device + * - Test to verify by Passing simple template struct + * - Test to verify by Passing simple explicit template struct + * - Test to verify by Passing struct with simple __device__ func() to struct memory + * - Test to verify by Passing struct which is initiazed out of order + * accessing same elements in order from device + * - Test to verify by Passing struct with bit fields operation + * accessing same elements in order from device + * - Test to verify by Passing the different hipLaunchParm options + + * Test source + * ------------------------ + * - catch/unit/kernel/hipLaunchParm.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipLaunchParm") { + hipMallocError = hipMalloc(reinterpret_cast(&result_d), + BLOCK_DIM_SIZE*sizeof(bool)); + hipHostMallocError = hipHostMalloc(reinterpret_cast(&result_h), + BLOCK_DIM_SIZE*sizeof(bool)); + hipMemsetError = hipMemset(result_d, false, BLOCK_DIM_SIZE); + + // Validating memory & initial value, for result_d, result_h + REQUIRE(hipMallocError == hipSuccess); + REQUIRE(hipHostMallocError == hipSuccess); + REQUIRE(hipMemsetError == hipSuccess); + + SECTION("check access from device") { + ResetValidationMem(); + hipLaunchKernelStruct_t1 hipLaunchKernelStruct_h1; + hipLaunchKernelStruct_h1.li = 1; + hipLaunchKernelStruct_h1.lf = 1.0; + hipLaunchKernelStruct_h1.result = false; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc1), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h1, + result_d); + ResultValidation(); + } + + SECTION("check padding") { + ResetValidationMem(); + hipLaunchKernelStruct_t2 hipLaunchKernelStruct_h2; + hipLaunchKernelStruct_h2.c1 = 'a'; + hipLaunchKernelStruct_h2.l1 = 1.0; + hipLaunchKernelStruct_h2.c2 = 'b'; + hipLaunchKernelStruct_h2.l2 = 2.0; + hipLaunchKernelStruct_h2.result = false; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc2), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h2, + result_d); + ResultValidation(); + } + + SECTION("check padding assigning int to char") { + ResetValidationMem(); + hipLaunchKernelStruct_t3 hipLaunchKernelStruct_h3; + hipLaunchKernelStruct_h3.bf1 = 1; + hipLaunchKernelStruct_h3.bf2 = 1; + hipLaunchKernelStruct_h3.l1 = 1.0; + hipLaunchKernelStruct_h3.bf3 = 1; + hipLaunchKernelStruct_h3.result = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc3), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h3, + result_d); + ResultValidation(); + } + + SECTION("Empty struct") { + ResetValidationMem(); + hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc4), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h4, + result_d); + ResultValidation(); + } + + SECTION("Passing struct with pointer object") { + ResetValidationMem(); + hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5; + char* cp_d5; // This is passed as pointer to struct member + // allocating memory for char pointer on device + HIP_CHECK(hipMalloc(reinterpret_cast(&cp_d5), sizeof(char))); + HIP_CHECK(hipMemset(cp_d5, 'p', sizeof(char))); + hipLaunchKernelStruct_h5.c1 = 'c'; + hipLaunchKernelStruct_h5.cp = cp_d5; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc5), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h5, + result_d); + ResultValidation(); + } + + SECTION("Passing struct with aligned(8)") { + ResetValidationMem(); + hipLaunchKernelStruct_t6 hipLaunchKernelStruct_h6; + hipLaunchKernelStruct_h6.c1 = 'c'; + hipLaunchKernelStruct_h6.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc6), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h6, + result_d); + // alignment is broken hence disabled the validation part + #if ENABLE_ALIGNMENT_TEST_SMALL_BAR + ResultValidation(); + #endif + } + + SECTION("Passing struct with aligned(16)") { + ResetValidationMem(); + hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7; + hipLaunchKernelStruct_h7.c1 = 'c'; + hipLaunchKernelStruct_h7.si = 1; + #if ENABLE_ALIGNMENT_TEST_SMALL_BAR // This is broken on small bar + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc7), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h7, + result_d); + ResultValidation(); + #endif + } + + SECTION("Passing struct with packed aligned to 4bytes") { + ResetValidationMem(); + hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8; + hipLaunchKernelStruct_h8.c1 = 'c'; + hipLaunchKernelStruct_h8.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8, + result_d); + // packed member broken on large and small bar setup. + #if ENABLE_PACKED_TEST + ResultValidation(); + #endif + } + + SECTION("Passing struct with packed to 4Bytes") { + ResetValidationMem(); + hipLaunchKernelStruct_t8A hipLaunchKernelStruct_h8A; + hipLaunchKernelStruct_h8A.c1 = 'c'; + hipLaunchKernelStruct_h8A.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8A), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8A, + result_d); + // packed member broken on large and small bar setup. + #if ENABLE_PACKED_TEST + ResultValidation(); + #endif + } + + SECTION("Passing struct with aligned(4) to 4Bytes") { + ResetValidationMem(); + hipLaunchKernelStruct_t8B hipLaunchKernelStruct_h8B; + hipLaunchKernelStruct_h8B.c1 = 'c'; + hipLaunchKernelStruct_h8B.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8B), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8B, + result_d); + // alignment is broken hence disabled the validation part + #if ENABLE_ALIGNMENT_TEST_SMALL_BAR + ResultValidation(); + #endif + } + + SECTION("Passing const struct object") { + ResetValidationMem(); + uint32_t* ip_d9; + // allocating memory for char pointer on device + HIP_CHECK(hipMalloc(reinterpret_cast(&ip_d9), sizeof(uint32_t))); + HIP_CHECK(hipMemset(ip_d9, 1, sizeof(uint32_t))); + // ip_d9 passed as pointer to struct member, struct.ip = &ip_d9 + const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_h9 = {'c', ip_d9}; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc9), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h9, + result_d); + #if ENABLE_DECLARE_INITIALIZATION_POINTER + ResultValidation(); + #endif + } + + SECTION("Passing struct with uintN_t") { + ResetValidationMem(); + hipLaunchKernelStruct_t10 hipLaunchKernelStruct_h10; + hipLaunchKernelStruct_h10.u64 = UINT64_MAX; + hipLaunchKernelStruct_h10.u32 = 1; + hipLaunchKernelStruct_h10.u8 = UINT8_MAX; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc10), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h10, + result_d); + ResultValidation(); + } + + SECTION("hipLaunchKernelStructFunc11") { + ResetValidationMem(); + hipLaunchKernelStruct_t11 hipLaunchKernelStruct_h11; + hipLaunchKernelStruct_h11.i1 = 1; + hipLaunchKernelStruct_h11.vint = 0; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc11), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h11, + result_d); + ResultValidation(); + } + + SECTION("Passing struct with simple class object") { + ResetValidationMem(); + hipLaunchKernelStruct_t12 hipLaunchKernelStruct_h12; + hipLaunchKernelStruct_h12.c1 = 'c'; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc12), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h12, + result_d); + #if ENABLE_CLASS_OBJ_ACCESS // access class obj from device broken + // Validation part of the struct, hipLaunchKernelStructFunc12 + ResultValidation(); + #endif + } + + SECTION("Passing struct with simple __device__ func()") { + ResetValidationMem(); + hipLaunchKernelStruct_t13 hipLaunchKernelStruct_h13; + hipLaunchKernelStruct_h13.i1 = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc13), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h13, + result_d); + ResultValidation(); + } + + SECTION("Passing struct with array variable") { + ResetValidationMem(); + hipLaunchKernelStruct_t14 hipLaunchKernelStruct_h14; + hipLaunchKernelStruct_h14.readint = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc14), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h14, + result_d); + ResultValidation(); + } + + SECTION("Passing struct with heap memory") { + ResetValidationMem(); + hipLaunchKernelStruct_t15 hipLaunchKernelStruct_h15; + hipLaunchKernelStruct_h15.c1 = 'c'; + + #if ENABLE_HEAP_MEMORY_ACCESS // causing page fault here, + // on small bar set + HIP_CHECK(hipMalloc(&hipLaunchKernelStruct_h15.heapmem, + BLOCK_DIM_SIZE*sizeof(int))); + HIP_CHECK(hipMemset(&hipLaunchKernelStruct_h15.heapmem, + 0, BLOCK_DIM_SIZE)); + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc15), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h15, + result_d); + ResultValidation(); + #endif + } + + SECTION("Passing simple template struct") { + ResetValidationMem(); + hipLaunchKernelStruct_t16 hipLaunchKernelStruct_h16; + hipLaunchKernelStruct_h16.t1 = 'c'; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc16), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h16, + result_d); + ResultValidation(); + } + + SECTION("Passing simple explicit template struct") { + ResetValidationMem(); + hipLaunchKernelStruct_t17 hipLaunchKernelStruct_h17; + hipLaunchKernelStruct_h17.t1 = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc17), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h17, + result_d); + ResultValidation(); + } + + SECTION("Passing struct with simple __device__ func()") { + ResetValidationMem(); + hipLaunchKernelStruct_t18 hipLaunchKernelStruct_h18; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc18), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h18, + result_d); + ResultValidation(); + } + + SECTION("Passing user defined stack") { + ResetValidationMem(); + hipLaunchKernelStruct_t19 hipLaunchKernelStruct_h19; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc19), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h19); + #if ENABLE_USER_STL + // Validation part of the struct, hipLaunchKernelStructFunc19 + HIPASSERT(hipLaunchKernelStruct_h19.stackSize() == 0); + #endif + } + + // Test: Passing struct which is initiazed out of order + // accessing same elements in order from device + SECTION("Passing struct which is initiazed out of order") { + ResetValidationMem(); + hipLaunchKernelStruct_t20 hipLaunchKernelStruct_h20; + hipLaunchKernelStruct_h20.name = 'A'; + hipLaunchKernelStruct_h20.rank = 2; + hipLaunchKernelStruct_h20.age = 42; + bool *result_d20, *result_h20; + #if ENABLE_OUT_OF_ORDER_INITIALIZATION + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc20), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h20, result_d); + ResultValidation(); + #endif + } + + SECTION("Passing struct with bit fields operation") { + ResetValidationMem(); + hipLaunchKernelStruct_t21 hipLaunchKernelStruct_h21 = + // out of order initalization + {2, 0}; + bool *result_d21, *result_h21; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc21), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h21, result_d); + #if ENABLE_BIT_FIELDS + ResultValidation(); + #endif + } + + SECTION("Passing the different hipLaunchParm options") { + float* Ad; + HIP_CHECK(hipMalloc(reinterpret_cast(&Ad), 1024)); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), size_t(1024), 1, 0, 0, Ad); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), 1024, dim3(1), 0, 0, Ad); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad); + + // Test: Passing macro to hipLaunchKernelGGL +#define KERNEL_CONFIG dim3(1024), dim3(1), 0, 0 + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), KERNEL_CONFIG, Ad); + + // Test: Same thing with templates: + int a; + float b; + hipLaunchKernelGGL(HIP_KERNEL_NAME(myKernel), + KERNEL_CONFIG, a, b); + +#define TYPE_PARAM_CONFIG int, float + hipLaunchKernelGGL(HIP_KERNEL_NAME(myKernel), + KERNEL_CONFIG, a, b); + + // Test: Passing hipLaunchKernelGGL inside another macro: + float e0; + MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024), + dim3(1), 0, 0, Ad), e0, j); + MY_LAUNCH_MACRO(WRAP(hipLaunchKernelGGL(vAdd, dim3(1024), + dim3(1), 0, 0, Ad)), e0, j); + +#ifdef EXTRA_PARENS_1 + // Don't wrap hipLaunchKernelGGL in extra set of parens: + MY_LAUNCH_MACRO((hipLaunchKernelGGL(vAdd, dim3(1024), + dim3(1), 0, 0, Ad)), e0, j); +#endif + + MY_LAUNCH(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1), + 0, 0, Ad), true, "firstCall"); + float* A; + float e1; + MY_LAUNCH_WITH_PAREN(hipMalloc(&A, 100), true, "launch2"); + +#ifdef EXTRA_PARENS_2 + // MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues. + MY_LAUNCH_WITH_PAREN(hipLaunchKernelGGL(vAdd, dim3(1024), + dim3(1), 0, 0, Ad), true, "firstCall"); +#endif + } + HIP_CHECK(hipHostFree(result_h)); + HIP_CHECK(hipFree(result_d)); +} diff --git a/projects/hip-tests/catch/unit/kernel/hipLaunchParmFunctor.cc b/projects/hip-tests/catch/unit/kernel/hipLaunchParmFunctor.cc new file mode 100644 index 0000000000..5b12ff52ae --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipLaunchParmFunctor.cc @@ -0,0 +1,459 @@ +/* +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 + +class HipFunctorTests { + public: + // Test that a class functor can be passed to hiplaunchparam + // and can be used in kernel + void TestForSimpleClassFunctor(void); + // Test that a templated class functor can be passed to hiplaunchparam + // and can be used in kernel + void TestForClassTemplateFunctor(void); + // Test that a class functor object ptr can be passed to hiplaunchparam + // and can be used in kernel + void TestForClassObjPtrFunctor(void); + // Test that a class object containing functor can be passed + // to hiplaunchparam and can be used in kernel + void TestForFunctorContainInClassObj(void); + // Test that a stuct functor can be passed to hiplaunchparam + // and can be used in kernel + void TestForSimpleStructFunctor(void); + // Test that a stuct functor object ptr can be passed to hiplaunchparam + // and can be used in kernel + void TestForStructObjPtrFunctor(void); + // Test that a templated struct functor can be passed to hiplaunchparam + // and can be used in kernel + void TestForStructTemplateFunctor(void); + // Test that a struct object containing functor can be + // passed to hiplaunchparam and can be used in kernel + void TestForFunctorContainInStructObj(void); +}; + +static const int BLOCK_DIM_SIZE = 1024; +static const int THREADS_PER_BLOCK = 1; + +// class functor tests + +// Simple doubler Functor +class DoublerFunctor{ + public: + __device__ int operator()(int x) { return x * 2;} +}; + +// simple doubler functor passed to kernel +__global__ void DoublerFunctorKernel( + DoublerFunctor doubler_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int result = doubler_(5); + deviceResult[x] = (result == 10); +} + +void HipFunctorTests::TestForSimpleClassFunctor(void) { + DoublerFunctor doubler; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(DoublerFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, doubler, deviceResults); + + // Validation part of TestForSimpleClassFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); +} + +// pointer functor passed to kernel +__global__ void PtrDoublerFunctorKernel( + DoublerFunctor *doubler_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int result = (*doubler_)(5); + deviceResult[x] = (result == 10); +} + +void HipFunctorTests::TestForClassObjPtrFunctor(void) { + DoublerFunctor* ptrdoubler = new DoublerFunctor[sizeof(int)]; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(PtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, ptrdoubler, deviceResults); + + // Validation part of TestForClassObjPtrFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); + delete[] ptrdoubler; +} + +class compare { + public: + template + __device__ bool operator()(const T1& v1, const T2& v2) { + return v1 > v2; + } +}; + +// template functor passed to kernel +__global__ void TemplateFunctorKernel( + compare compare_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + deviceResult[x] = compare_(2.2, 2.1); + deviceResult[x] = compare_(2, 1); + deviceResult[x] = compare_('b', 'a'); +} + +void HipFunctorTests::TestForClassTemplateFunctor(void) { + compare comparefunctor; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(TemplateFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults); + + // Validation part of TestForClassTemplateFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); +} + + +// Doubler calculator +class DoublerCalculator { + public: + int a, result; + // fucntor contained in class object + DoublerFunctor doubler; +}; + +// doubler functor conatined in class obj passed to kernel +__global__ void DoublerCalculatorFunctorKernel( + DoublerCalculator doubler_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int result = doubler_.doubler(doubler_.a); + deviceResult[x] = (doubler_.result == result); +} + +void HipFunctorTests::TestForFunctorContainInClassObj(void) { + DoublerCalculator Doubler; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + Doubler.a = 5; + Doubler.result = 10; + // pass comparefunctor to hipLaunchParm + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults); + + // Validation part of TestForStructTemplateFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); +} + +// Struct functor tests + +// Simple doubler Functor +struct sDoublerFunctor { + public: + __device__ int operator()(int x) { return x * 2;} +}; + + +// simple sturct doubler functor passed to kernel +__global__ void structDoublerFunctorKernel( + sDoublerFunctor doubler_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int result = doubler_(5); + deviceResult[x] = (result == 10); +} + +void HipFunctorTests::TestForSimpleStructFunctor(void) { + sDoublerFunctor doubler; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(structDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, doubler, deviceResults); + + // Validation part of TestForSimpleStructFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); +} + +// ptr functor passed to kernel +__global__ void structPtrDoublerFunctorKernel( + sDoublerFunctor *doubler_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int result = (*doubler_)(5); + deviceResult[x] = (result == 10); +} + +void HipFunctorTests::TestForStructObjPtrFunctor(void) { + sDoublerFunctor* ptrdoubler = new sDoublerFunctor[sizeof(int)]; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(structPtrDoublerFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, ptrdoubler, deviceResults); + + // Validation part of TestForStructObjPtrFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); + delete[] ptrdoubler; +} + +struct sCompare { + public: + template< typename T1, typename T2 > + __device__ bool operator()(const T1& v1, const T2& v2) { + return v1 > v2; + } +}; + +// template functor passed to kernel +__global__ void structTemplateFunctorKernel( + sCompare compare_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + deviceResult[x] = compare_(2.2, 2.1); + deviceResult[x] = compare_(2, 1); + deviceResult[x] = compare_('b', 'a'); +} + +void HipFunctorTests::TestForStructTemplateFunctor(void) { + sCompare comparefunctor; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + + // pass comparefunctor to hipLaunchKernelGGL + hipLaunchKernelGGL(structTemplateFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, comparefunctor, deviceResults); + + // Validation part of TestForStructTemplateFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); +} + +// Doubler calculator struct +struct sDoublerCalculator { + public: + int a, result; + // fucntor contained in class object + DoublerFunctor doubler; +}; + + + +// doubler functor contained in struct passed to kernel +__global__ void DoublerCalculatorFunctorKernel( + sDoublerCalculator doubler_, + bool* deviceResult) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int result = doubler_.doubler(doubler_.a); + deviceResult[x] = (doubler_.result == result); +} + +void HipFunctorTests::TestForFunctorContainInStructObj(void) { + sDoublerCalculator Doubler; + bool *deviceResults, *hostResults; + HIP_CHECK(hipMalloc(&deviceResults, BLOCK_DIM_SIZE*sizeof(bool))); + HIP_CHECK(hipHostMalloc(&hostResults, BLOCK_DIM_SIZE*sizeof(bool))); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + // initialize to false, will be set to + // true if the functor is called in device code + hostResults[k] = false; + } + + Doubler.a = 5; + Doubler.result = 10; + HIP_CHECK(hipMemcpy(deviceResults, hostResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyHostToDevice)); + + + // pass comparefunctor to hipLaunchKernelGGL + hipLaunchKernelGGL(DoublerCalculatorFunctorKernel, dim3(BLOCK_DIM_SIZE), + dim3(THREADS_PER_BLOCK), 0, 0, Doubler, deviceResults); + + // Validation part of TestForStructTemplateFunctor + HIP_CHECK(hipMemcpy(hostResults, deviceResults, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost)); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + REQUIRE(hostResults[k] == true); + HIP_CHECK(hipHostFree(hostResults)); + HIP_CHECK(hipFree(deviceResults)); +} + +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup KernelTest +* `void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args)` - +* Method to invocate kernel functions +*/ + +/** + * Test Description + * ------------------------ + * - Test that a class functor can be passed to hiplaunchparam + * and can be used in kernel. + * - Test that a templated class functor can be passed to hiplaunchparam + * and can be used in kernel. + * - Test that a class functor object ptr can be passed to hiplaunchparam + * and can be used in kernel. + * - Test that a class object containing functor can be passed to hiplaunchparam + * and can be used in kernel + * - Test that a stuct functor can be passed to hiplaunchparam + * and can be used in kernel + * - Test that a stuct functor object ptr can be passed to hiplaunchparam + * and can be used in kernel + * - Test that a templated struct functor can be passed to hiplaunchparam + * and can be used in kernel + * - Test that a struct object containing functor can be passed to hiplaunchparam + * and can be used in kernel + + * Test source + * ------------------------ + * - catch/unit/kernel/hipLaunchParmFunctor.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ + +TEST_CASE("Unit_hipLaunchParmFunctor") { + HipFunctorTests FunctorTests; + + SECTION("test for simple class functor") { + FunctorTests.TestForSimpleClassFunctor(); + } + SECTION("test for class objptr functor") { + FunctorTests.TestForClassObjPtrFunctor(); + } + SECTION("test for class templete functor") { + FunctorTests.TestForClassTemplateFunctor(); + } + SECTION("test for simple struct functor") { + FunctorTests.TestForSimpleStructFunctor(); + } + SECTION("test for struct objptr functor") { + FunctorTests.TestForStructObjPtrFunctor(); + } + SECTION("test for struct templete functor") { + FunctorTests.TestForStructTemplateFunctor(); + } + SECTION("test for functor contain in classobj") { + FunctorTests.TestForFunctorContainInClassObj(); + } + SECTION("test for functor contain in structobj") { + FunctorTests.TestForFunctorContainInStructObj(); + } +}