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

[Jan Vesely via llvm-commits]

jvesely 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:
----------------
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


Repository:
  rL LLVM

http://reviews.llvm.org/D20299