[PATCH] D20299: AMDGPU: Export target workitem related builtins

[Tom Stellard via llvm-commits]

tstellarAMD added inline comments.

================
Comment at: lib/CodeGen/CGBuiltin.cpp:7506-7519
@@ -7492,1 +7505,16 @@
+  // amdgcn workitem
+  case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
+    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0, 256);
+  case AMDGPU::BI__builtin_amdgcn_workitem_id_y:
+    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_y, 0, 256);
+  case AMDGPU::BI__builtin_amdgcn_workitem_id_z:
+    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 256);
+
+  // r600 workitem
+  case AMDGPU::BI__builtin_r600_read_tidig_x:
+    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_x, 0, 256);
+  case AMDGPU::BI__builtin_r600_read_tidig_y:
+    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_y, 0, 256);
+  case AMDGPU::BI__builtin_r600_read_tidig_z:
+    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0, 256);
   default:
----------------
jvesely wrote:
> tstellarAMD wrote:
> > Are you sure 256 is the upper bounds for these?
> I'm pretty sure it's not. There was a short discussion earlier in this revision.
> OpenGL requires at least 1024x1024x64 (1024 total) for compute shaders, so I'd say hw supports at least those sizes.
> EG/CM ISA specs don't say.
> SI/CI/VI ISA specs say 1024.
> Mesa exposes either 256 or 2048.
> 
> Larger sets can be faked using GDS, but since there is no lower bound in OpenCL it'd be nice to have (efficient) hw limits here
I think we should use 1024 here,  Using a number too low will generate incorrect code.  It's also probably same to assume the 1024 is limit for EG/CM too.

A follow up improvement might be to check for OpenCL related function attributes, like reqd_work_group_size and use that to emit a smaller range.


Repository:
  rL LLVM

http://reviews.llvm.org/D20299