Templated Kernels

tests/runtime/CMakeLists.txt

@@ -25,6 +25,30 @@
 
 set(CHIP_SKIP_TEST 77)
 
+add_custom_target(TestTemplatedKernelArgs ALL
+  COMMAND ${CMAKE_BINARY_DIR}/bin/hipcc.bin -c 
+    ${CMAKE_CURRENT_SOURCE_DIR}/inputs/TemplatedKernelAddOne.hip 
+    ${CMAKE_CURRENT_SOURCE_DIR}/inputs/TemplatedKernelAddTwo.hip
+    &&
+    ${CMAKE_BINARY_DIR}/bin/hipcc.bin -o TestTemplatedKernelArgs
+    ${CMAKE_CURRENT_SOURCE_DIR}/TestTemplatedKernels.hip
+    ${CMAKE_CURRENT_BINARY_DIR}/TemplatedKernelAddOne.o 
+    ${CMAKE_CURRENT_BINARY_DIR}/TemplatedKernelAddTwo.o
+
+  DEPENDS 
+    ${CMAKE_CURRENT_SOURCE_DIR}/inputs/TemplatedKernelAddOne.hip  
+    ${CMAKE_CURRENT_SOURCE_DIR}/inputs/TemplatedKernelAddTwo.hip
+    ${CMAKE_CURRENT_SOURCE_DIR}/TestTemplatedKernels.hip
+    hipcc.bin CHIP
+
+  COMMENT "Building TestTemplatedKernelArgs"
+)
+add_test(NAME TestTemplatedKernelArgs COMMAND TestTemplatedKernelArgs)
+set_tests_properties(TestTemplatedKernelArgs PROPERTIES
+         FAIL_REGULAR_EXPRESSION "device function is already registered"
+         TIMEOUT 60
+         )
+
 function(add_hip_runtime_test MAIN_SOURCE)
   get_filename_component(EXEC_NAME ${MAIN_SOURCE} NAME_WLE)
   set_source_files_properties(${MAIN_SOURCE} PROPERTIES LANGUAGE CXX)

tests/runtime/TestTemplatedKernels.hip

@@ -0,0 +1,24 @@
+#include <hip/hip_runtime.h>
+
+/**
+ * Both of these funcitons rely on a templated kernelAddOne<int> function
+ * launchKernelAddTwo launches the kernel twice. 
+ * Similar pattern is commonly seen in hip-tests
+ * 
+ * If the templated kernel is replaced by a regular function we get a compilation error:
+ * TemplatedKernelAddTwo.hip:(.text+0x0): multiple definition of `__device_stub__kernelAddOne(int*)';
+*/
+void launchKernelAddOne(int* data);
+void launchKernelAddTwo(int* data);
+
+int main() {
+  int* data;
+  hipHostMalloc((void**)&data, sizeof(int));
+  data[0] = 0;
+  assert(data[0] == 0);
+  launchKernelAddOne(data);
+  launchKernelAddTwo(data);
+
+  hipFree(data);
+  return 0;
+}

tests/runtime/inputs/TemplatedKernelAddOne.hip

@@ -0,0 +1,10 @@
+#include <hip/hip_runtime.h>
+
+template<typename T>
+__global__ void kernelAddOne(T* data) {
+    data[0] += 1;
+}
+
+void launchKernelAddOne(int* data) {
+  hipLaunchKernelGGL(kernelAddOne<int>, dim3(1), dim3(1), 0, 0, data);
+}

tests/runtime/inputs/TemplatedKernelAddTwo.hip

@@ -0,0 +1,11 @@
+#include <hip/hip_runtime.h>
+
+template<typename T>
+__global__ void kernelAddOne(T* data) {
+    data[0] += 1;
+}
+
+void launchKernelAddTwo(int* data) {
+  hipLaunchKernelGGL(kernelAddOne<int>, dim3(1), dim3(1), 0, 0, data);
+  hipLaunchKernelGGL(kernelAddOne<int>, dim3(1), dim3(1), 0, 0, data);
+}