Merge pull request #490 from jdmaia/hip-guard-acceleratorfor2dNB

[HIP] Including kernel launch parameter guard on accelerator_for2dNB
Including guard on accelerator_for2dNB against invalid kernel configurations if GRID_HIP
2026-07-30 07:23:28 +01:00 · 2026-05-06 14:51:30 -04:00 · 2026-05-06 13:26:33 -05:00
1 changed files with 14 additions and 16 deletions
@@ -432,6 +432,7 @@ accelerator_inline int acceleratorSIMTlane(int Nsimd) {

 #define accelerator_for2dNB( iter1, num1, iter2, num2, nsimd, ... )	\
  {									\
+    if (num1*num2) { \
      typedef uint64_t Iterator;						\
      auto lambda = [=] accelerator					\
        (Iterator iter1,Iterator iter2,Iterator lane ) mutable {		\
@@ -441,13 +442,10 @@ accelerator_inline int acceleratorSIMTlane(int Nsimd) {
      dim3 hip_threads(nsimd, nt, 1);					 \
      dim3 hip_blocks ((num1+nt-1)/nt,num2,1); \
      if(hip_threads.x * hip_threads.y * hip_threads.z <= 64){ \
-      hipLaunchKernelGGL(LambdaApply64,hip_blocks,hip_threads,		\
-   	                 0,computeStream,						\
-			 num1,num2,nsimd, lambda);			\
+        LambdaApply64<<<hip_blocks,hip_threads,0,computeStream>>>(num1,num2,nsimd,lambda);			\
      } else { \
-      hipLaunchKernelGGL(LambdaApply,hip_blocks,hip_threads,		\
-			 0,computeStream,				\
-			 num1,num2,nsimd, lambda);			\
+        LambdaApply<<<hip_blocks,hip_threads,0,computeStream>>>(num1,num2,nsimd,lambda);			\
+      } \
    } \
  }
Author	SHA1	Message	Date
Peter BoyleandGitHub	e2d607f6c7	Merge pull request #490 from jdmaia/hip-guard-acceleratorfor2dNB [HIP] Including kernel launch parameter guard on accelerator_for2dNB	2026-05-06 14:51:30 -04:00
Julio Maia	66da4e0657	Including guard on accelerator_for2dNB against invalid kernel configurations if GRID_HIP	2026-05-06 13:26:33 -05:00