[mlir][transform] Fail if thre is no `thread_dim_mapping`

[guraypp]

Summary

map_nested_foreach_to_threads has implicit loop mapping; for details, see below. While it's correct, it does not make use of warp programming.

for(i)    --> threadIdx.x
 for(j)   --> threadIdx.y
  for(k)  --> threadIdx.z

In certain circumstances, the compiler can automatically determine the optimal loop mapping, but not always. We anticipate explicit mapping from the user until we find a fully performant solution. So this revision adds failiure if there is no thread_dim_mapping.

mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp

static DiagnosedSilenceableFailure rewriteOneForeachThreadToGpuThreads(
   if (foreachThreadOp.getNumResults() > 0)
     return failureHelper(
         "only bufferized scf.foreach_thread lowers to gpu.thread_id");
 
   if (foreachThreadOp.getNumThreads().size() > 3)
     return failureHelper(
         "scf.foreach_thread with rank > 3 does not lower to gpu.thread_id");
 
+  if (foreachThreadOp.getThreadDimMapping().empty())
+    return failureHelper(
+        "scf.foreach_thread is missing thread_dim_mapping attribute.");
+
   auto potentialBlockDim = foreachThreadOp.getPermutedNumThreads(rewriter);
   if (failed(potentialBlockDim) ||
       llvm::any_of(*potentialBlockDim, [](OpFoldResult ofr) {
         return !getConstantIntValue(ofr).has_value();
       })) {
     return failureHelper("unsupported dynamic blockdim size");
   }
 

mlir/test/Dialect/GPU/transform-gpu-failing.mlir

 ^bb0(%arg0: !pdl.operation):  
   %funcop = transform.structured.match ops{["func.func"]} in %arg0   
   // expected-error @below {{Trying to launch a GPU kernel with gridDim = (65535, 65535, 1) blockDim = (1, 1, 1). It is larger than the limits.}}
   %1 = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch }
 }
 
 // -----
 
+
+// -----
+
+!type = memref<2 x 32 x f32>
+!type1d = memref<32 x f32>
+
+// CHECK-LABEL: func.func @saxpy2d_no_barrier(
+func.func @no_thread_map(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream : !gpu.async.token) -> !type {
+  %one = arith.constant 1 : index
+  %c12 = arith.constant 12 : index
+  %c9 = arith.constant 9 : index
+  %c7 = arith.constant 7 : index
+  %name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
+            threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
+  {
+    scf.foreach_thread (%i, %j) in (%c7, %c9) {
+        %4 = memref.load %x[%i, %j] : !type
+        %5 = memref.load %y[%i, %j] : !type
+        %6 = math.fma %alpha, %4, %5 : f32
+        memref.store %6, %y[%i, %j] : !type
+     }  
+    gpu.terminator
+  }
+  return %y : !type
+}
+
+transform.with_pdl_patterns {
+^bb0(%arg0: !pdl.operation):
+  transform.sequence %arg0 failures(propagate) {
+  ^bb1(%arg1: !pdl.operation):
+    %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0     
+    // expected-error @below {{scf.foreach_thread is missing thread_dim_mapping attribute.}}   
+    transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1] }
+  }
+}
 No newline at end of file