EXSWHTEC-105 - Implement tests for hipModuleOccupancyMaxActiveBlocksPerMultiprocessor and hipModuleOccupancyMaxPotentialBlockSize APIs #47

Change-Id: I4e51bb999f43c4f304ed4712dadc3fc1aae3f7a4


[ROCm/hip-tests commit: bb791d8e86]

projects/hip-tests/catch/hipTestMain/config/config_amd_linux

@@ -54,6 +54,8 @@
         "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
         "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported",
         "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters",
+        "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters",
+        "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters",
         "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic",
         "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic",
         "Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Positive_Basic",

projects/hip-tests/catch/hipTestMain/config/config_amd_windows

@@ -119,6 +119,8 @@
         "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
         "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported",
         "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters",
+        "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters",
+        "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters",
         "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic",
         "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic",
         "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Positive_Basic",

projects/hip-tests/catch/unit/occupancy/CMakeLists.txt

@@ -4,9 +4,21 @@ set(TEST_SRC
   hipOccupancyMaxActiveBlocksPerMultiprocessor_old.cc
   hipOccupancyMaxPotentialBlockSize.cc
   hipOccupancyMaxPotentialBlockSize_old.cc
+  hipModuleOccupancyMaxPotentialBlockSize.cc
+  hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc
+  hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc
+  hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc
   hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.cc
 )
 
+add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/simple_kernel.code
+                   COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/simple_kernel.cc -o simple_kernel.code
+                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/simple_kernel.cc)
+
+add_custom_target(simple_kernel ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/simple_kernel.code)
+
 hip_add_exe_to_target(NAME OccupancyTest
                       TEST_SRC ${TEST_SRC}
                       TEST_TARGET_NAME build_tests)
+
+add_dependencies(OccupancyTest simple_kernel)

projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc

@@ -0,0 +1,92 @@
+/*
+Copyright (c) 2022 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.
+*/
+/*
+Testcase Scenarios :
+Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation - Test correct
+execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessor for diffrent parameter values
+Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters - Test unsuccessful
+execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessor api when parameters are invalid
+*/
+#include "occupancy_common.hh"
+
+TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters") {
+  hipModule_t module;
+  hipFunction_t function;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  // Get potential blocksize
+  HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0));
+
+  // Common negative tests
+  MaxActiveBlocksPerMultiprocessorNegative(
+      [&function](int* numBlocks, int blockSize, size_t dynSharedMemPerBlk) {
+        return hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, function, blockSize,
+                                                                  dynSharedMemPerBlk);
+      },
+      blockSize);
+
+  HIP_CHECK(hipModuleUnload(module));
+}
+
+TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation") {
+  hipDeviceProp_t devProp;
+  hipModule_t module;
+  hipFunction_t function;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  SECTION("dynSharedMemPerBlk = 0") {
+    // Get potential blocksize
+    HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0));
+
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize, &function](int* numBlocks) {
+          return hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, function, blockSize,
+                                                                    0);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
+  SECTION("dynSharedMemPerBlk = sharedMemPerBlock") {
+    // Get potential blocksize
+    HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function,
+                                                      devProp.sharedMemPerBlock, 0));
+
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize, devProp, &function](int* numBlocks) {
+          return hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, function, blockSize,
+                                                                    devProp.sharedMemPerBlock);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}

projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc

@@ -0,0 +1,103 @@
+/*
+Copyright (c) 2022 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.
+*/
+/*
+Testcase Scenarios :
+Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation - Test
+correct execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags for diffrent
+parameter values
+Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters - Test
+unsuccessful execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags api when
+parameters are invalid
+*/
+#include "occupancy_common.hh"
+
+TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters") {
+  hipModule_t module;
+  hipFunction_t function;
+  int numBlocks = 0;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  // Get potential blocksize
+  HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0));
+
+  // Common negative tests
+  MaxActiveBlocksPerMultiprocessorNegative(
+      [&function](int* numBlocks, int blockSize, size_t dynSharedMemPerBlk) {
+        return hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+            numBlocks, function, blockSize, dynSharedMemPerBlk, hipOccupancyDefault);
+      },
+      blockSize);
+
+  SECTION("Flag is invalid") {
+    // Only default flag is supported
+    HIP_CHECK_ERROR(hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+                        &numBlocks, function, blockSize, 0, 2),
+                    hipErrorInvalidValue);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}
+
+TEST_CASE(
+    "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation") {
+  hipDeviceProp_t devProp;
+  hipModule_t module;
+  hipFunction_t function;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  SECTION("dynSharedMemPerBlk = 0") {
+    // Get potential blocksize
+    HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0));
+
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize, &function](int* numBlocks) {
+          return hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+              numBlocks, function, blockSize, 0, hipOccupancyDefault);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
+  SECTION("dynSharedMemPerBlk = sharedMemPerBlock") {
+    // Get potential blocksize
+    HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function,
+                                                      devProp.sharedMemPerBlock, 0));
+
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize, devProp, &function](int* numBlocks) {
+          return hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+              numBlocks, function, blockSize, devProp.sharedMemPerBlock, hipOccupancyDefault);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}

projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc

@@ -0,0 +1,75 @@
+/*
+Copyright (c) 2022 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.
+*/
+/*
+Testcase Scenarios :
+Unit_hipModuleOccupancyMaxPotentialBlockSize_Positive_RangeValidation - Test correct execution of
+hipModuleOccupancyMaxPotentialBlockSize for diffrent parameter values
+Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters - Test unsuccessful execution of
+hipModuleOccupancyMaxPotentialBlockSize api when parameters are invalid
+*/
+#include "occupancy_common.hh"
+
+TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters") {
+  hipModule_t module;
+  hipFunction_t function;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  // Common negative tests
+  MaxPotentialBlockSizeNegative([&function](int* gridSize, int* blockSize) {
+    return hipModuleOccupancyMaxPotentialBlockSize(gridSize, blockSize, function, 0, 0);
+  });
+
+  HIP_CHECK(hipModuleUnload(module));
+}
+
+TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Positive_RangeValidation") {
+  hipDeviceProp_t devProp;
+  hipModule_t module;
+  hipFunction_t function;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  SECTION("dynSharedMemPerBlk = 0, blockSizeLimit = 0") {
+    MaxPotentialBlockSize(
+        [&function](int* gridSize, int* blockSize) {
+          return hipModuleOccupancyMaxPotentialBlockSize(gridSize, blockSize, function, 0, 0);
+        },
+        devProp.maxThreadsPerBlock);
+  }
+
+  SECTION("dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock") {
+    MaxPotentialBlockSize(
+        [&function, devProp](int* gridSize, int* blockSize) {
+          return hipModuleOccupancyMaxPotentialBlockSize(
+              gridSize, blockSize, function, devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock);
+        },
+        devProp.maxThreadsPerBlock);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}

projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc

@@ -0,0 +1,87 @@
+/*
+Copyright (c) 2022 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.
+*/
+/*
+Testcase Scenarios :
+Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Positive_RangeValidation - Test correct
+execution of hipModuleOccupancyMaxPotentialBlockSizeWithFlags for diffrent parameter values
+Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters - Test unsuccessful
+execution of hipModuleOccupancyMaxPotentialBlockSizeWithFlags api when parameters are invalid
+*/
+#include "occupancy_common.hh"
+
+TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters") {
+  hipModule_t module;
+  hipFunction_t function;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  // Common negative tests
+  MaxPotentialBlockSizeNegative([&function](int* gridSize, int* blockSize) {
+    return hipModuleOccupancyMaxPotentialBlockSizeWithFlags(gridSize, blockSize, function, 0, 0,
+                                                            hipOccupancyDefault);
+  });
+
+  SECTION("Flag is invalid") {
+    // Only default flag is supported
+    HIP_CHECK_ERROR(
+        hipModuleOccupancyMaxPotentialBlockSizeWithFlags(&gridSize, &blockSize, function, 0, 0, 2),
+        hipErrorInvalidValue);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}
+
+TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Positive_RangeValidation") {
+  hipDeviceProp_t devProp;
+  hipModule_t module;
+  hipFunction_t function;
+
+  HIP_CHECK(hipFree(nullptr));
+
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  SECTION("dynSharedMemPerBlk = 0, blockSizeLimit = 0") {
+    MaxPotentialBlockSize(
+        [&function](int* gridSize, int* blockSize) {
+          return hipModuleOccupancyMaxPotentialBlockSizeWithFlags(gridSize, blockSize, function, 0,
+                                                                  0, hipOccupancyDefault);
+        },
+        devProp.maxThreadsPerBlock);
+  }
+
+  SECTION("dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock") {
+    MaxPotentialBlockSize(
+        [&function, devProp](int* gridSize, int* blockSize) {
+          return hipModuleOccupancyMaxPotentialBlockSizeWithFlags(
+              gridSize, blockSize, function, devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock,
+              hipOccupancyDefault);
+        },
+        devProp.maxThreadsPerBlock);
+  }
+
+  HIP_CHECK(hipModuleUnload(module));
+}

projects/hip-tests/catch/unit/occupancy/occupancy_common.hh

@@ -66,7 +66,5 @@ template <typename F> void MaxActiveBlocksPerMultiprocessorNegative(F func, int
   SECTION("numBlocks is nullptr") {
     HIP_CHECK_ERROR(func(nullptr, blockSize, 0), hipErrorInvalidValue);
   }
-  SECTION("Block size is 0") { 
-    HIP_CHECK_ERROR(func(&numBlocks, 0, 0), hipErrorInvalidValue);
-  }
+  SECTION("Block size is 0") { HIP_CHECK_ERROR(func(&numBlocks, 0, 0), hipErrorInvalidValue); }
 }

projects/hip-tests/catch/unit/occupancy/simple_kernel.cc

@@ -0,0 +1,25 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+
+extern "C" __global__ void SimpleKernel(int* a, int* b) {
+    int tx = threadIdx.x;
+    b[tx] = a[tx];
+}