From c1f57da0738bcc148c85406c80e1ed3e91410e55 Mon Sep 17 00:00:00 2001 From: Rupam Chetia Date: Fri, 12 Mar 2021 17:09:04 +0530 Subject: [PATCH] SWDEV-238517 - [dtest] Adding additional scenarios for block and grid dimension. 1. Corner case and negative test scenarios added for block and grid dimension. Change-Id: I094faf02570fec101f688462712934b94ceb37e1 [ROCm/hip commit: b1d78e4096ac9b6639139387171363a7f8fa7a8c] --- .../module/hipExtModuleLaunchKernel.cpp | 206 ++++++++++++++- .../module/hipModuleLaunchKernel.cpp | 243 +++++++++++++++++- .../tests/src/runtimeApi/module/matmul.cpp | 4 +- 3 files changed, 444 insertions(+), 9 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp index 55cf0fba57..ea03677ea8 100755 --- a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -1,5 +1,5 @@ /* - Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2021 - present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -23,6 +23,12 @@ without concurrency flag and verify the time difference between them 4. hipExtModuleLaunchKernel API verifying the kernel execution time of a particular kernel. 5. hipExtModuleLaunchKernel API verifying the kernel execution time by disabling the time flag + 6. hipExtModuleLaunchKernel API verifying Corner Scenarios for Grid and Block dimensions + 7. hipModuleLaunchKernel Work Group tests => + - (block.x * block.y * block.z) <= Work Group Size + where block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ + - (block.x * block.y * block.z) > Work Group Size + where block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ Scenarios 2 and 3 concurrency verification scenarios are not included in HIT command as firmware currently does not support the concurrency in the same stream based on the flag. @@ -33,10 +39,11 @@ * TEST_NAMED: %t hipExtModuleLaunchKernel_NegativeTests --tests 1 EXCLUDE_HIP_PLATFORM nvidia * TEST_NAMED: %t hipExtModuleLaunchKernel_KernelExecutionTime --tests 4 EXCLUDE_HIP_PLATFORM nvidia * TEST_NAMED: %t hipExtModuleLaunchKernel_DisabledEventTimeFlag --tests 5 EXCLUDE_HIP_PLATFORM nvidia + * TEST_NAMED: %t hipExtModuleLaunchKernel_CornerScenarios --tests 6 EXCLUDE_HIP_PLATFORM nvidia + * TEST_NAMED: %t hipExtModuleLaunchKernel_WorkGroup --tests 7 EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ - -#include +#include #include "test_common.h" #include "hip/hip_ext.h" @@ -47,9 +54,18 @@ #define FourSec "FourSecKernel" #define TwoSec "TwoSecKernel" #define globalDevVar "deviceGlobal" +#define dummyKernel "dummyKernel" #define FOURSEC_KERNEL 4999 #define TWOSEC_KERNEL 2999 +struct gridblockDim { + unsigned int gridX; + unsigned int gridY; + unsigned int gridZ; + unsigned int blockX; + unsigned int blockY; + unsigned int blockZ; +}; class ModuleLaunchKernel { int N = 64; int SIZE = N*N; @@ -61,7 +77,7 @@ class ModuleLaunchKernel { hipModule_t Module; hipDeviceptr_t deviceGlobal; hipFunction_t MultKernel, SixteenSecKernel, FourSecKernel, - TwoSecKernel, KernelandExtraParamKernel; + TwoSecKernel, KernelandExtraParamKernel, DummyKernel; struct { int clockRate; void* _Ad; @@ -69,8 +85,11 @@ class ModuleLaunchKernel { void* _Cd; int _n; } args1, args2; + struct { + } args3; size_t size1; size_t size2; + size_t size3; size_t deviceGlobalSize; public : void AllocateMemory(); @@ -78,6 +97,8 @@ class ModuleLaunchKernel { void ModuleLoad(); bool Module_Negative_tests(); bool ExtModule_Negative_tests(); + bool ExtModule_Corner_tests(); + bool Module_WorkGroup_Test(); bool ExtModule_KernelExecutionTime(); bool ExtModule_ConcurencyCheck_GlobalVar(int conc_flag); bool ExtModule_ConcurrencyCheck_TimeVer(); @@ -116,6 +137,7 @@ void ModuleLaunchKernel::AllocateMemory() { args2.clockRate = clkRate; size1 = sizeof(args1); size2 = sizeof(args2); + size3 = sizeof(args3); HIPCHECK(hipEventCreate(&start_event1)); HIPCHECK(hipEventCreate(&end_event1)); HIPCHECK(hipEventCreate(&start_event2)); @@ -134,6 +156,7 @@ void ModuleLaunchKernel::ModuleLoad() { Module, KernelandExtra)); HIPCHECK(hipModuleGetFunction(&FourSecKernel, Module, FourSec)); HIPCHECK(hipModuleGetFunction(&TwoSecKernel, Module, TwoSec)); + HIPCHECK(hipModuleGetFunction(&DummyKernel, Module, dummyKernel)); HIPCHECK(hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, globalDevVar)); } @@ -355,6 +378,54 @@ bool ModuleLaunchKernel::ExtModule_Negative_tests() { printf("hipExtModuleLaunchKernel failed for max values to block dimension"); testStatus = false; } + // Passing 0 as value for all dimensions + err = hipExtModuleLaunchKernel(MultKernel, 0, 0, 0, + 0, + 0, + 0, 0, + stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for 0 as value for all dimensions"); + testStatus = false; + } + // Passing 0 as value for x dimension + err = hipExtModuleLaunchKernel(MultKernel, 0, 1, 1, + 0, + 1, + 1, 0, + stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for 0 as value for x dimension"); + testStatus = false; + } + // Passing 0 as value for y dimension + err = hipExtModuleLaunchKernel(MultKernel, 1, 0, 1, + 1, + 0, + 1, 0, + stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for 0 as value for y dimension"); + testStatus = false; + } + // Passing 0 as value for z dimension + err = hipExtModuleLaunchKernel(MultKernel, 1, 1, 0, + 1, + 1, + 0, 0, + stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for 0 as value for z dimension"); + testStatus = false; + } // Passing both kernel and extra params err = hipExtModuleLaunchKernel(KernelandExtraParamKernel, 1, 1, 1, 1, 1, 1, 0, stream1, reinterpret_cast(¶ms), @@ -378,6 +449,40 @@ bool ModuleLaunchKernel::ExtModule_Negative_tests() { printf("hipExtModuleLaunchKernel failed for max group size"); testStatus = false; } + // Block dimension X = Max Allowed + 1 + err = hipExtModuleLaunchKernel(MultKernel, 1, 1, 1, + deviceProp.maxThreadsDim[0]+1, + 1, + 1, 0, stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for (MaxBlockDimX + 1)"); + testStatus = false; + } + // Block dimension Y = Max Allowed + 1 + err = hipExtModuleLaunchKernel(MultKernel, 1, 1, 1, + 1, + deviceProp.maxThreadsDim[1]+1, + 1, 0, stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for (MaxBlockDimY + 1)"); + testStatus = false; + } + // Block dimension Z = Max Allowed + 1 + err = hipExtModuleLaunchKernel(MultKernel, 1, 1, 1, + 1, + 1, + deviceProp.maxThreadsDim[2]+1, 0, stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed for (MaxBlockDimZ + 1)"); + testStatus = false; + } + // Passing invalid config data in extra params void *config3[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1, @@ -393,6 +498,95 @@ bool ModuleLaunchKernel::ExtModule_Negative_tests() { return testStatus; } +bool ModuleLaunchKernel::ExtModule_Corner_tests() { + bool testStatus = true; + HIPCHECK(hipSetDevice(0)); + hipError_t err; + AllocateMemory(); + ModuleLoad(); + void *config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args3, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size3, + HIP_LAUNCH_PARAM_END}; + hipDeviceProp_t deviceProp; + hipGetDeviceProperties(&deviceProp, 0); + unsigned int maxblockX = deviceProp.maxThreadsDim[0]; + unsigned int maxblockY = deviceProp.maxThreadsDim[1]; + unsigned int maxblockZ = deviceProp.maxThreadsDim[2]; + struct gridblockDim test[6] = {{1, 1, 1, maxblockX, 1, 1}, + {1, 1, 1, 1, maxblockY, 1}, + {1, 1, 1, 1, 1, maxblockZ}, + {UINT32_MAX, 1, 1, 1, 1, 1}, + {1, UINT32_MAX, 1, 1, 1, 1}, + {1, 1, UINT32_MAX, 1, 1, 1}}; + + for (int i = 0; i < 6; i++) { + err = hipExtModuleLaunchKernel(DummyKernel, + test[i].gridX, + test[i].gridY, + test[i].gridZ, + test[i].blockX, + test[i].blockY, + test[i].blockZ, + 0, + stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err != hipSuccess) { + printf("hipExtModuleLaunchKernel failed (%u, %u, %u) and (%u, %u, %u)", + test[i].gridX, test[i].gridY, test[i].gridZ, + test[i].blockX, test[i].blockY, test[i].blockZ); + testStatus = false; + } + } + DeAllocateMemory(); + return testStatus; +} + +bool ModuleLaunchKernel::Module_WorkGroup_Test() { + bool testStatus = true; + HIPCHECK(hipSetDevice(0)); + hipError_t err; + AllocateMemory(); + ModuleLoad(); + void *config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args3, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size3, + HIP_LAUNCH_PARAM_END}; + hipDeviceProp_t deviceProp; + hipGetDeviceProperties(&deviceProp, 0); + double cuberootVal = + cbrt(static_cast(deviceProp.maxThreadsPerBlock)); + uint32_t cuberoot_floor = floor(cuberootVal); + uint32_t cuberoot_ceil = ceil(cuberootVal); + // Scenario: (block.x * block.y * block.z) <= Work Group Size where + // block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ + err = hipExtModuleLaunchKernel(DummyKernel, + 1, 1, 1, + cuberoot_floor, cuberoot_floor, cuberoot_floor, + 0, stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err != hipSuccess) { + printf("hipExtModuleLaunchKernel failed block dimensions (%u, %u, %u)", + cuberoot_floor, cuberoot_floor, cuberoot_floor); + testStatus = false; + } + // Scenario: (block.x * block.y * block.z) > Work Group Size where + // block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ + err = hipExtModuleLaunchKernel(DummyKernel, + 1, 1, 1, + cuberoot_ceil, cuberoot_ceil, cuberoot_ceil + 1, + 0, stream1, NULL, + reinterpret_cast(&config1), + nullptr, nullptr, 0); + if (err == hipSuccess) { + printf("hipExtModuleLaunchKernel failed block dimensions (%u, %u, %u)", + cuberoot_ceil, cuberoot_ceil, cuberoot_ceil); + testStatus = false; + } + DeAllocateMemory(); + return testStatus; +} + int main(int argc, char* argv[]) { bool testStatus = true; HipTest::parseStandardArguments(argc, argv, false); @@ -408,6 +602,10 @@ int main(int argc, char* argv[]) { testStatus &= kernelLaunch.ExtModule_KernelExecutionTime(); } else if (p_tests == 5) { testStatus &= kernelLaunch.ExtModule_Disabled_Timingflag(); + } else if (p_tests == 6) { + testStatus &= kernelLaunch.ExtModule_Corner_tests(); + } else if (p_tests == 7) { + testStatus &= kernelLaunch.Module_WorkGroup_Test(); } else { failed("Didnt receive any valid option.\n"); } diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp index 2e72dcb6a3..40cbccd2d0 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp @@ -1,5 +1,5 @@ /* - Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2021 - present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -18,15 +18,23 @@ */ /* Test Scenarios 1. hipModuleLaunchKernel Negative Scenarios + 2. hipModuleLaunchKernel Corner Scenarios for Grid and Block dimensions + 3. hipModuleLaunchKernel Work Group tests => + - (block.x * block.y * block.z) <= Work Group Size + where block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ + - (block.x * block.y * block.z) > Work Group Size + where block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ */ /* HIT_START * BUILD_CMD: matmul.code %hc --genco %S/matmul.cpp -o matmul.code * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 - * TEST: %t + * TEST: %t --tests 0x1 + * TEST: %t --tests 0x2 + * TEST: %t --tests 0x3 * HIT_END */ - +#include #include "test_common.h" #define fileName "matmul.code" @@ -35,7 +43,16 @@ #define KernelandExtra "KernelandExtraParams" #define FourSec "FourSecKernel" #define TwoSec "TwoSecKernel" +#define dummyKernel "dummyKernel" +struct gridblockDim { + unsigned int gridX; + unsigned int gridY; + unsigned int gridZ; + unsigned int blockX; + unsigned int blockY; + unsigned int blockZ; +}; bool Module_Negative_tests() { bool testStatus = true; @@ -93,6 +110,50 @@ bool Module_Negative_tests() { printf("hipModuleLaunchKernel failed for max values to block dimension"); testStatus = false; } + // Passing 0 as value for all dimensions + err = hipModuleLaunchKernel(MultKernel, 0, 0, 0, + 0, + 0, + 0, 0, + stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for 0 as value for all dimensions"); + testStatus = false; + } + // Passing 0 as value for x dimension + err = hipModuleLaunchKernel(MultKernel, 0, 1, 1, + 0, + 1, + 1, 0, + stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for 0 as value for x dimension"); + testStatus = false; + } + // Passing 0 as value for y dimension + err = hipModuleLaunchKernel(MultKernel, 1, 0, 1, + 1, + 0, + 1, 0, + stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for 0 as value for y dimension"); + testStatus = false; + } + // Passing 0 as value for z dimension + err = hipModuleLaunchKernel(MultKernel, 1, 1, 0, + 1, + 1, + 0, 0, + stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for 0 as value for z dimension"); + testStatus = false; + } // Passing both kernel and extra params err = hipModuleLaunchKernel(KernelandExtraParamKernel, 1, 1, 1, 1, 1, 1, 0, stream1, @@ -114,6 +175,36 @@ bool Module_Negative_tests() { printf("hipModuleLaunchKernel failed for max group size"); testStatus = false; } + // Block dimension X = Max Allowed + 1 + err = hipModuleLaunchKernel(MultKernel, 1, 1, 1, + deviceProp.maxThreadsDim[0]+1, + 1, + 1, 0, stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for (MaxBlockDimX + 1)"); + testStatus = false; + } + // Block dimension Y = Max Allowed + 1 + err = hipModuleLaunchKernel(MultKernel, 1, 1, 1, + 1, + deviceProp.maxThreadsDim[1]+1, + 1, 0, stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for (MaxBlockDimY + 1)"); + testStatus = false; + } + // Block dimension Z = Max Allowed + 1 + err = hipModuleLaunchKernel(MultKernel, 1, 1, 1, + 1, + 1, + deviceProp.maxThreadsDim[2]+1, 0, stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed for (MaxBlockDimZ + 1)"); + testStatus = false; + } // Passing invalid config data to extra params void *config3[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1, @@ -132,9 +223,153 @@ bool Module_Negative_tests() { return testStatus; } +bool Module_GridBlock_Corner_Tests() { + bool testStatus = true; + HIPCHECK(hipSetDevice(0)); + hipError_t err; + struct { + } args1; + hipFunction_t DummyKernel; + size_t size1; + size1 = sizeof(args1); + hipModule_t Module; + hipStream_t stream1; + hipDeviceptr_t *Ad; + hipDevice_t device; +#ifdef __HIP_PLATFORM_NVCC__ + hipCtx_t context; + HIPCHECK(hipDeviceGet(&device, 0)); + HIPCHECK(hipCtxCreate(&context, 0, device)); +#endif + HIPCHECK(hipModuleLoad(&Module, fileName)); + HIPCHECK(hipModuleGetFunction(&DummyKernel, Module, dummyKernel)); + void *config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args1, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1, + HIP_LAUNCH_PARAM_END}; + HIPCHECK(hipStreamCreate(&stream1)); + // Passing Max int value to block dimensions + hipDeviceProp_t deviceProp; + hipGetDeviceProperties(&deviceProp, 0); + unsigned int maxblockX = deviceProp.maxThreadsDim[0]; + unsigned int maxblockY = deviceProp.maxThreadsDim[1]; + unsigned int maxblockZ = deviceProp.maxThreadsDim[2]; +#ifdef __HIP_PLATFORM_NVCC__ + unsigned int maxgridX = deviceProp.maxGridSize[0]; + unsigned int maxgridY = deviceProp.maxGridSize[1]; + unsigned int maxgridZ = deviceProp.maxGridSize[2]; +#else + unsigned int maxgridX = UINT32_MAX; + unsigned int maxgridY = UINT32_MAX; + unsigned int maxgridZ = UINT32_MAX; +#endif + struct gridblockDim test[6] = {{1, 1, 1, maxblockX, 1, 1}, + {1, 1, 1, 1, maxblockY, 1}, + {1, 1, 1, 1, 1, maxblockZ}, + {maxgridX, 1, 1, 1, 1, 1}, + {1, maxgridY, 1, 1, 1, 1}, + {1, 1, maxgridZ, 1, 1, 1}}; + for (int i = 0; i < 6; i++) { + err = hipModuleLaunchKernel(DummyKernel, + test[i].gridX, + test[i].gridY, + test[i].gridZ, + test[i].blockX, + test[i].blockY, + test[i].blockZ, + 0, + stream1, NULL, + reinterpret_cast(&config1)); + if (err != hipSuccess) { + printf("hipModuleLaunchKernel failed (%u, %u, %u) and (%u, %u, %u)", + test[i].gridX, test[i].gridY, test[i].gridZ, + test[i].blockX, test[i].blockY, test[i].blockZ); + testStatus = false; + } + } + HIPCHECK(hipStreamDestroy(stream1)); + HIPCHECK(hipModuleUnload(Module)); +#ifdef __HIP_PLATFORM_NVCC__ + hipCtxDestroy(context); +#endif + return testStatus; +} + +bool Module_WorkGroup_Test() { + bool testStatus = true; + HIPCHECK(hipSetDevice(0)); + hipError_t err; + struct { + } args1; + hipFunction_t DummyKernel; + size_t size1; + size1 = sizeof(args1); + hipModule_t Module; + hipStream_t stream1; + hipDeviceptr_t *Ad; + hipDevice_t device; +#ifdef __HIP_PLATFORM_NVCC__ + hipCtx_t context; + HIPCHECK(hipDeviceGet(&device, 0)); + HIPCHECK(hipCtxCreate(&context, 0, device)); +#endif + HIPCHECK(hipModuleLoad(&Module, fileName)); + HIPCHECK(hipModuleGetFunction(&DummyKernel, Module, dummyKernel)); + void *config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args1, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1, + HIP_LAUNCH_PARAM_END}; + HIPCHECK(hipStreamCreate(&stream1)); + // Passing Max int value to block dimensions + hipDeviceProp_t deviceProp; + hipGetDeviceProperties(&deviceProp, 0); + double cuberootVal = + cbrt(static_cast(deviceProp.maxThreadsPerBlock)); + uint32_t cuberoot_floor = floor(cuberootVal); + uint32_t cuberoot_ceil = ceil(cuberootVal); + // Scenario: (block.x * block.y * block.z) <= Work Group Size where + // block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ + err = hipModuleLaunchKernel(DummyKernel, + 1, 1, 1, + cuberoot_floor, cuberoot_floor, cuberoot_floor, + 0, stream1, NULL, + reinterpret_cast(&config1)); + if (err != hipSuccess) { + printf("hipModuleLaunchKernel failed block dimensions (%u, %u, %u)", + cuberoot_floor, cuberoot_floor, cuberoot_floor); + testStatus = false; + } + // Scenario: (block.x * block.y * block.z) > Work Group Size where + // block.x < MaxBlockDimX , block.y < MaxBlockDimY and block.z < MaxBlockDimZ + err = hipModuleLaunchKernel(DummyKernel, + 1, 1, 1, + cuberoot_ceil, cuberoot_ceil, cuberoot_ceil + 1, + 0, stream1, NULL, + reinterpret_cast(&config1)); + if (err == hipSuccess) { + printf("hipModuleLaunchKernel failed block dimensions (%u, %u, %u)", + cuberoot_ceil, cuberoot_ceil, cuberoot_ceil); + testStatus = false; + } + HIPCHECK(hipStreamDestroy(stream1)); + HIPCHECK(hipModuleUnload(Module)); +#ifdef __HIP_PLATFORM_NVCC__ + hipCtxDestroy(context); +#endif + return testStatus; +} + int main(int argc, char* argv[]) { bool testStatus = true; - testStatus = Module_Negative_tests(); + HipTest::parseStandardArguments(argc, argv, true); + if (p_tests == 0x1) { + testStatus = Module_Negative_tests(); + } else if (p_tests == 0x2) { + testStatus = Module_GridBlock_Corner_Tests(); + } else if (p_tests == 0x3) { + testStatus = Module_WorkGroup_Test(); + } else { + printf("Invalid Test Case \n"); + exit(1); + } if (testStatus) { passed(); } else { diff --git a/projects/hip/tests/src/runtimeApi/module/matmul.cpp b/projects/hip/tests/src/runtimeApi/module/matmul.cpp index 2d1641443e..8f803e8ffd 100755 --- a/projects/hip/tests/src/runtimeApi/module/matmul.cpp +++ b/projects/hip/tests/src/runtimeApi/module/matmul.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021 - present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -79,3 +79,5 @@ extern "C" __global__ void FourSecKernel(int clockrate) { } } +extern "C" __global__ void dummyKernel() { +}