<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:x="urn:schemas-microsoft-com:office:excel" xmlns:p="urn:schemas-microsoft-com:office:powerpoint" xmlns:a="urn:schemas-microsoft-com:office:access" xmlns:dt="uuid:C2F41010-65B3-11d1-A29F-00AA00C14882" xmlns:s="uuid:BDC6E3F0-6DA3-11d1-A2A3-00AA00C14882" xmlns:rs="urn:schemas-microsoft-com:rowset" xmlns:z="#RowsetSchema" xmlns:b="urn:schemas-microsoft-com:office:publisher" xmlns:ss="urn:schemas-microsoft-com:office:spreadsheet" xmlns:c="urn:schemas-microsoft-com:office:component:spreadsheet" xmlns:odc="urn:schemas-microsoft-com:office:odc" xmlns:oa="urn:schemas-microsoft-com:office:activation" xmlns:html="http://www.w3.org/TR/REC-html40" xmlns:q="http://schemas.xmlsoap.org/soap/envelope/" xmlns:rtc="http://microsoft.com/officenet/conferencing" xmlns:D="DAV:" xmlns:Repl="http://schemas.microsoft.com/repl/" xmlns:mt="http://schemas.microsoft.com/sharepoint/soap/meetings/" xmlns:x2="http://schemas.microsoft.com/office/excel/2003/xml" xmlns:ppda="http://www.passport.com/NameSpace.xsd" xmlns:ois="http://schemas.microsoft.com/sharepoint/soap/ois/" xmlns:dir="http://schemas.microsoft.com/sharepoint/soap/directory/" xmlns:ds="http://www.w3.org/2000/09/xmldsig#" xmlns:dsp="http://schemas.microsoft.com/sharepoint/dsp" xmlns:udc="http://schemas.microsoft.com/data/udc" xmlns:xsd="http://www.w3.org/2001/XMLSchema" xmlns:sub="http://schemas.microsoft.com/sharepoint/soap/2002/1/alerts/" xmlns:ec="http://www.w3.org/2001/04/xmlenc#" xmlns:sp="http://schemas.microsoft.com/sharepoint/" xmlns:sps="http://schemas.microsoft.com/sharepoint/soap/" xmlns:xsi="http://www.w3.org/2001/XMLSchema-instance" xmlns:udcs="http://schemas.microsoft.com/data/udc/soap" xmlns:udcxf="http://schemas.microsoft.com/data/udc/xmlfile" xmlns:udcp2p="http://schemas.microsoft.com/data/udc/parttopart" xmlns:wf="http://schemas.microsoft.com/sharepoint/soap/workflow/" xmlns:dsss="http://schemas.microsoft.com/office/2006/digsig-setup" xmlns:dssi="http://schemas.microsoft.com/office/2006/digsig" xmlns:mdssi="http://schemas.openxmlformats.org/package/2006/digital-signature" xmlns:mver="http://schemas.openxmlformats.org/markup-compatibility/2006" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns:mrels="http://schemas.openxmlformats.org/package/2006/relationships" xmlns:spwp="http://microsoft.com/sharepoint/webpartpages" xmlns:ex12t="http://schemas.microsoft.com/exchange/services/2006/types" xmlns:ex12m="http://schemas.microsoft.com/exchange/services/2006/messages" xmlns:pptsl="http://schemas.microsoft.com/sharepoint/soap/SlideLibrary/" xmlns:spsl="http://microsoft.com/webservices/SharePointPortalServer/PublishedLinksService" xmlns:Z="urn:schemas-microsoft-com:" xmlns:st="" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv=Content-Type content="text/html; charset=us-ascii">
<meta name=Generator content="Microsoft Word 12 (filtered medium)">
<style>
<!--
/* Font Definitions */
@font-face
{font-family:"Cambria Math";
panose-1:2 4 5 3 5 4 6 3 2 4;}
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri","sans-serif";}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
span.EmailStyle17
{mso-style-type:personal-compose;
font-family:"Calibri","sans-serif";
color:windowtext;}
.MsoChpDefault
{mso-style-type:export-only;}
@page Section1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.Section1
{page:Section1;}
-->
</style>
<!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang=EN-US link=blue vlink=purple>
<div class=Section1>
<p class=MsoNormal>Is there a current way to specify that an instruction or
function call cannot be duplicated and thus any optimizations that might want
to duplicate this instruction would fail?<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>The problem deals with barrier in OpenCL 1.0. One of the
conditions of using barrier is that if a barrier exists inside of control flow,
every thread in a work-group must execute the barrier instruction(6.11.9). <o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>However, in this simple CL code:<br>
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics: enable<o:p></o:p></p>
<p class=MsoNormal>#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics:
enable<o:p></o:p></p>
<p class=MsoNormal>__kernel void <o:p></o:p></p>
<p class=MsoNormal>KMeansMapReduceAtomic(const int num_attributes,<o:p></o:p></p>
<p class=MsoNormal>
const int num_objects,<o:p></o:p></p>
<p class=MsoNormal>
__global int*
delta_d<o:p></o:p></p>
<p class=MsoNormal>
) <o:p></o:p></p>
<p class=MsoNormal>{<o:p></o:p></p>
<p class=MsoNormal> __local
int clusterCount[256];<o:p></o:p></p>
<p class=MsoNormal> __local
int sTemp[1]; // amd opencl needed this to be an array<o:p></o:p></p>
<p class=MsoNormal> const unsigned int point_id =
get_local_id(0);<o:p></o:p></p>
<p class=MsoNormal> int index = 0;<o:p></o:p></p>
<p class=MsoNormal> int i, addr;<o:p></o:p></p>
<p class=MsoNormal> int xx = get_local_id(0);<o:p></o:p></p>
<p class=MsoNormal> clusterCount[xx] = 0;<o:p></o:p></p>
<p class=MsoNormal> if(get_local_id(0) == 0){<o:p></o:p></p>
<p class=MsoNormal> sTemp[0] = 0; //sTemp is for prefix
sum<o:p></o:p></p>
<p class=MsoNormal> }<o:p></o:p></p>
<p class=MsoNormal> barrier(CLK_LOCAL_MEM_FENCE);<o:p></o:p></p>
<p class=MsoNormal> int idWithinCluster = 300; // anthing other
then zero<o:p></o:p></p>
<p class=MsoNormal> if (point_id < num_objects) {<o:p></o:p></p>
<p class=MsoNormal> idWithinCluster =
atom_add(&clusterCount<o:p></o:p></p>
<p class=MsoNormal>
[index],1);<o:p></o:p></p>
<p class=MsoNormal> }<o:p></o:p></p>
<p class=MsoNormal>
barrier(CLK_LOCAL_MEM_FENCE);<o:p></o:p></p>
<p class=MsoNormal style='text-indent:.5in'>int numMembers = 2; <o:p></o:p></p>
<p class=MsoNormal> if(idWithinCluster
== 0) {<o:p></o:p></p>
<p class=MsoNormal> clusterCount[index]
= atom_add(&sTemp[0], numMembers);//This holds the prefix offset<o:p></o:p></p>
<p class=MsoNormal> }<o:p></o:p></p>
<p class=MsoNormal style='text-indent:.5in'> delta_d[xx] =
clusterCount[index];<o:p></o:p></p>
<p class=MsoNormal>}<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>produces bitcode file which has 3 barriers.<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>The problem is now that the second if/barrier pair:<o:p></o:p></p>
<p class=MsoNormal>if (point_id < num_objects) {<o:p></o:p></p>
<p class=MsoNormal> idWithinCluster =
atom_add(&clusterCount<o:p></o:p></p>
<p class=MsoNormal>
[index],1);<o:p></o:p></p>
<p class=MsoNormal> }<o:p></o:p></p>
<p class=MsoNormal>
barrier(CLK_LOCAL_MEM_FENCE);<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>is transformed into flow control equivalent to :<o:p></o:p></p>
<p class=MsoNormal>if (point_id >= num_objects) {<o:p></o:p></p>
<p class=MsoNormal> barrier(CLK_LOCAL_MEM_FENCE);<o:p></o:p></p>
<p class=MsoNormal>} else {<o:p></o:p></p>
<p class=MsoNormal> idWithinCluster =
atom_add(&clusterCount<o:p></o:p></p>
<p class=MsoNormal>
[index],1);<o:p></o:p></p>
<p class=MsoNormal> barrier(CLK_LOCAL_MEM_FENCE);<o:p></o:p></p>
<p class=MsoNormal>}<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>which violates opencl, which can cause undefined behavior on
the underlying hardware, as each barrier is unique. <o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>So we want to disable all optimizations around barrier
instructions, but not in other cases when no barrier instruction exists. One
way to do this is to mark an instruction as not being copyable, but is there a
method of doing this in LLVM?<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>Also, this barrier does not map to llvm.barrier because
llvm.barrier only seems to worry about memory operations and not synchronization
between threads.<o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal>Thanks for any help,<o:p></o:p></p>
<p class=MsoNormal>Micah<o:p></o:p></p>
<p class=MsoNormal> <o:p></o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
<p class=MsoNormal><o:p> </o:p></p>
</div>
</body>
</html>