public void Run()
{
Console.WriteLine("================== " + GetType().Name + "================== ");
string source = @"
__kernel void myOpenClFunc ( __global float* cl_input, __global float* cl_output )
{
size_t i = get_global_id(0);
cl_output[i] = cl_input[i] + cl_input[i];
}; ";
/************ Initialize OpenClWithGCN ***************************************/
OpenClWithGCN gprog = new OpenClWithGCN();
OpenClEnvironment env = gprog.env;
bool success = gprog.GcnCompile(source);
Console.Write(env.lastMessage);
if (!success)
{
Assert.Fail();
return;
}
/************ Create some random data *******************************************/
// create some random data for testing
var random = new Random();
const int count = 1024 * 1024;
const int dataSz = count * sizeof(float);
float[] data = (from i in Enumerable.Range(0, count) select(float) random.NextDouble()).ToArray();
/************ Build and run the kernel *******************************************/
Kernel kernel = env.program.CreateKernel("myOpenClFunc");
Mem cl_input = env.context.CreateBuffer(MemoryFlags.ReadOnly, dataSz);
Mem cl_output = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);
env.cmdQueue.EnqueueWriteBuffer(cl_input, true, 0, dataSz, data);
kernel.Arguments[0].SetValue(cl_input);
kernel.Arguments[1].SetValue(cl_output);
env.cmdQueue.EnqueueNDRangeKernel(kernel, count, 256);
env.cmdQueue.Finish();
/************ Read back and Validate the results ***********************************/
float[] results = new float[count];
env.cmdQueue.EnqueueReadBufferAndWait(cl_output, results, dataSz);
int correct = Enumerable.Range(0, count).Where(i => results[i] == data[i] * 2).Count();
Console.WriteLine("{0} - Computed {1}/{2} correct values!",
correct == count ? "PASS" : "FAIL", correct.ToString(), count.ToString());
Assert.AreEqual(correct, count);
}
private void Form1_Load(object sender, EventArgs e)
{
string gpuInfo = OpenClWithGCN.CheckGPUAndVersion();
if (gpuInfo.Length > 0)
txtOutput.AppendText(gpuInfo);
if (gcn.env == null)
txtOutput.AppendText("Unable to find an AMD platform.");
/////////////////// Copy needed exe/dll files to output folder ///////////////////
if (!Directory.Exists(outputFolder))
Directory.CreateDirectory(outputFolder);
//if (!File.Exists(outputFolder + "\\OpenClWithGCN.dll"))
File.Copy("OpenClWithGCN.dll", outputFolder + "\\OpenClWithGCN.dll", true);
//if (!File.Exists(outputFolder + "\\Asm4GCN.exe"))
File.Copy("Asm4GCN.exe", outputFolder + "\\Asm4GCN.exe", true);
//if (!File.Exists(outputFolder + "\\NOpenCL.dll"))
File.Copy("NOpenCL.dll", outputFolder + "\\NOpenCL.dll", true);
/////////////////// Setup AutoCompletes for GCN asm window //////////////////////
autocompleteMenu1.MaximumSize = new System.Drawing.Size(700, 3000);
var columnWidth = new int[] { 125, 575 };
foreach (var i in GcnTools.ISA_DATA.ISA_Insts)
{
int img = i.name[0]=='s'?0:i.name[0]=='v'?1:i.name[0]=='d'?2:i.name[0]=='b'?3:4;
autocompleteMenu1.AddItem(new AC.MulticolumnAutocompleteItem(new[] { i.name, i.iSANotes },
i.name) { ColumnWidth = columnWidth, ImageIndex = img });
}
autocompleteMenu1.AddItem(new AC.MulticolumnAutocompleteItem(new[] {
"__asm4GCN", "Begins a Asm4GCN Block" }
, "__asm4GCN") { ColumnWidth = columnWidth, ImageIndex = 4 }
);
autocompleteMenu1.AddItem(new AC.MulticolumnAutocompleteItem(new[] {
"__kernel", "Begins an OpenCL Block" },
"__kernel") { ColumnWidth = columnWidth, ImageIndex = 4 }
);
changeSinceLastSave = false;
this.Text = "Welcome to " + Application.ProductName;
}
public void Run()
{
Console.WriteLine("================== " + GetType().Name + "================== ");
string source = @"
__asm4GCN myAsmFunc ( float*, float*)
{
// ======== Pre-Loaded Registers =========
// s[2:3] - PTR_UAV_TABLE
// s[2:3] +0x60 - base_resource_const1(#T)
// s[2:3] +0x68 - base_resource_const2(#T)
// s[4:7] - IMM_CONST_BUFFER0
// s[4:7] +0x00 - Grid Size
// s[4:7] +0x04 - Local Size
// s[4:7] +0x18 - baseGlobalId
// s[8:11] -IMM_CONST_BUFFER1
// s[8:11]+0x00 - param1 offset
// s[8:11]+0x04 - param2 offset
// s[8:11]+0x08 - param3 offset
// s1 - threadgroupId
// s12 - groupId
// v0 - laneId
// Use a #define to shorten tbuffer_load instruction
#define _F32_ 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT]
// Variable decelerations with specific registers (pre-loaded regs)
s8u resConsts s[2:3] // base resource constants
s16u uavBuff s[4:7], paramsOffset s[8:11], groupId s12
v4b laneId v0 // v0 is pre-loaded with lane id
// Setup Resource Constant
s_load_dwordx4 s16u resConst1, resConsts, 0x68
// establish GlobalID
s_buffer_load_dword s4u baseGlobalId,uavBuff, 0x18
s_buffer_load_dword s4u localSize, uavBuff, 0x04
s_waitcnt lgkmcnt(0)
v_mov_b32 v4u vLocalSize, localSize
v_mul_i32_i24 vLocalSize, groupId, vLocalSize
v_add_i32 v4u localSizeIdx, vcc, laneId, vLocalSize
v_add_i32 v4u vGlobalID, vcc, baseGlobalId, localSizeIdx
v_lshlrev_b32 v4u vGlobalOffset, 2, vGlobalID
// Fetch the value that we want to multiply by 2.
s_buffer_load_dword s4u sPara1Ptr, paramsOffset, 0x00
s_load_dwordx4 s16u resConst0, resConsts, 0x60
s_waitcnt lgkmcnt(0)
v_add_i32 v4b baseOffset, vcc, sPara1Ptr, vGlobalOffset
tbuffer_load_format_x v4b val, baseOffset, resConst0, _F32_
s_waitcnt vmcnt(0)
// Perform the 'sum = val * 2'
v_mul_f32 v4f sum, 2.0, val
// Write the results back to memory.
s_buffer_load_dword s4u param2Offset, paramsOffset, 0x04
s_waitcnt lgkmcnt(0)
s_waitcnt vmcnt(0)
v_add_i32 v4u dstOffset, vcc, param2Offset, vGlobalOffset
free param2Offset // var is freed and register returned to pool.
tbuffer_store_format_x sum, dstOffset, resConst1, _F32_
// Exit the kernel
s_endpgm
}";
/************ Initialize OpenClWithGCN ***************************************/
OpenClWithGCN gprog = new OpenClWithGCN();
OpenClEnvironment env = gprog.env;
bool success = gprog.GcnCompile(source);
Console.Write(env.lastMessage);
if (!success)
{
Assert.Fail();
return;
}
/************ Create some random data *******************************************/
// create some random data for testing
var random = new Random();
const int count = 1024 * 1024;
const int dataSz = count * sizeof(float);
float[] data = (from i in Enumerable.Range(0, count) select(float) random.NextDouble()).ToArray();
/************ Build and run the kernel *******************************************/
Kernel kernel = env.program.CreateKernel("myAsmFunc");
Mem cl_input = env.context.CreateBuffer(MemoryFlags.ReadOnly, dataSz);
Mem cl_output = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);
env.cmdQueue.EnqueueWriteBuffer(cl_input, true, 0, dataSz, data);
kernel.Arguments[0].SetValue(cl_input);
kernel.Arguments[1].SetValue(cl_output);
env.cmdQueue.EnqueueNDRangeKernel(kernel, count, 256);
env.cmdQueue.Finish();
/************ Read back and Validate the results ***********************************/
float[] results = new float[count];
env.cmdQueue.EnqueueReadBufferAndWait(cl_output, results, dataSz);
int correct = Enumerable.Range(0, count).Where(i => results[i] == data[i] * 2).Count();
Console.WriteLine("{0} - Computed {1}/{2} correct values!",
correct == count ? "PASS" : "FAIL", correct.ToString(), count.ToString());
Assert.AreEqual(correct, count);
}
public void Run()
{
// The example below is an OpenCL example by Derek Gerstmann(UWA). It's been modified for NOpenCL use.
Console.WriteLine("================== " + GetType().Name + "================== ");
/************ Create and build a program from our OpenCL-C source code ***************/
// const string source = GCN_NS.Code.DevCode;
string source = @"
__asm4GCN myAsmFunc ( float*, float*)
{
#define _ZeroInHex_ 0x00
#define _32Float_ 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT]
// Loc| Binary |BinaryExt
s_buffer_load_dword s0, s[4:7], 0x04 // 000|C2000504|
s_buffer_load_dword s1, s[4:7], 0x18 // 004|C2008518|
s_waitcnt lgkmcnt(0) // 008|BF8C007F|
s_min_u32 s0, s0, 0x0000ffff // 00C|8380FF00|0000FFFF
s_buffer_load_dword s4, s[8:11], _ZeroInHex_ // 014|C2020900|
v_mov_b32 v1, s0 // 018|7E020200|
v_mul_i32_i24 v1, s12, v1 // 01C|1202020C|
v_add_i32 v0, vcc, v0, v1 // 020|4A000300|
v_add_i32 v0, vcc, s1, v0 // 024|4A000001|
v_lshlrev_b32 v0, 2, v0 // 028|34000082|
s_load_dwordx4 s[12:15], s[2:3], 0x60 // 02C|C0860360|
s_waitcnt lgkmcnt(0) // 030|BF8C007F|
v_add_i32 v1, vcc, s4, v0 // 034|4A020004|
tbuffer_load_format_x v1, v1, s[12:15], _32Float_ // 038|EBA01000|80030101
s_buffer_load_dword s0, s[8:11], 0x04 // 040|C2000904|
s_load_dwordx4 s[4:7], s[2:3], 0x68 // 044|C0820368|
s_waitcnt lgkmcnt(0) // 048|BF8C007F|
v_add_i32 v0, vcc, s0, v0 // 04C|4A000000|
s_waitcnt vmcnt(0) // 050|BF8C0F70|
v_add_f32 v1, v1, v1 // 054|10020301|
tbuffer_store_format_x v1, v0, s[4:7], _32Float_ // 058|EBA41000|80010100
s_endpgm // 060|BF810000|
};
// This is not used here but is what generates above when de-assembled.
// __kernel and __asm4GCN blocks can be used in the same clProgram
__kernel void myOpenClFunc ( __global float* cl_input, __global float* cl_output )
{
size_t i = get_global_id(0);
cl_output[i] = cl_input[i] + cl_input[i];
}; ";
/************ Initialize OpenClWithGCN ***************************************/
// OpenClEnvironment env = SetupOpenClEnvironment(); // for manual setup
OpenClWithGCN gprog = new OpenClWithGCN();
OpenClEnvironment env = gprog.env; // Let just use the default environment
bool success = gprog.GcnCompile(source);
Console.Write(env.lastMessage);
if (!success)
{
Assert.Fail();
return;
}
/************ Create some random data *******************************************/
// create some random data for testing
var random = new Random();
const int count = 1024 * 1024;
const int dataSz = count * sizeof(float);
float[] data = (from i in Enumerable.Range(0, count) select(float) random.NextDouble()).ToArray();
/************ Create a kernel from our modProgram *******************************/
Kernel kernel = env.program.CreateKernel("myAsmFunc");
/************ Allocate cl_input, and fill with data ********************************/
// OpenCL: cl_mem cl_input = clCreateBuffer(context, CL_MEM_READ_ONLY, dataSz, host_ptr, NULL);
Mem cl_input = env.context.CreateBuffer(MemoryFlags.ReadOnly, dataSz);
/************ Create an cl_output memory buffer for our results *****************/
// OpenCL: cl_mem cl_output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSz, host_ptr, NULL);
Mem cl_output = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);
/************ Copy our host buffer of random values to cl_input device buffer ******/
// OpenCL: clEnqueueWriteBuffer(cmdQueue, cl_input, CL_TRUE, 0, dataSz, data, 0, NULL, NULL);
env.cmdQueue.EnqueueWriteBuffer(cl_input, true, 0, dataSz, data);
/************ Set the arguments to our kernel, and enqueue it for execution ********/
// OpenCL: clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_input);
kernel.Arguments[0].SetValue(cl_input);
kernel.Arguments[1].SetValue(cl_output);
/************ Enqueue and run the kernel *******************************************/
// OpenCL: clEnqueueNDRangeKernel(cmdQueue, kernel, 1, NULL, &globalMem &localMem, 0, NULL,NULL);
env.cmdQueue.EnqueueNDRangeKernel(kernel, count, 256);
/************ Force command queue to get processed, wait until all commands finish **/
env.cmdQueue.Finish();
/************ Read back the results ************************************************/
// OpenCL: clEnqueueReadBuffer(cmdQueue, cl_output, CL_TRUE, 0, dataSz, results, 0, NULL, NULL);
float[] results = new float[count];
env.cmdQueue.EnqueueReadBufferAndWait(cl_output, results, dataSz);
/************ Validate our results *************************************************/
int correct = 0;
for (int i = 0; i < count; i++)
{
correct += (results[i] == data[i] + data[i]) ? 1 : 0;
}
// int correct = Enumerable.Range(0,count).Where(i=>results[i]==data[i]*2).Count();
/************ Print a brief summary detailing the results **************************/
Console.WriteLine("{0} - Computed {1}/{2} correct values!",
correct == count ? "PASS" : "FAIL", correct.ToString(), count.ToString());
Assert.AreEqual(correct, count);
}
public void Run()
{
Console.WriteLine("================== " + GetType().Name + "================== ");
// Create and build a program from our OpenCL-C source code
string source = @"__asm4GCN myAsmFunc (float*,float*)
{
// ======== Pre-Loaded Registers =========
// s[2:3] - UAV table Pointer
// s[2:3] +0x60 - base_resource_const1(#T)
// s[2:3] +0x68 - base_resource_const2(#T)
// s[4:7] - Imm Const Buffer 0
// s[4:7] +0x00 - Grid Size
// s[4:7] +0x04 - Local Size
// s[4:7] +0x18 - baseGlobalId
// s[8:11] - Imm Const Buffer 1
// s[8:11]+0x00 - param1 offset
// s[8:11]+0x04 - param2 offset
// s[8:11]+0x08 - param3 offset (not used here)
// s12 - Group ID
// v0 - Local ID
s_mov_b32 m0, 0x00010000
s_movk_i32 s0, 0x0000
s_movk_i32 s1, 0x0000
s_movk_i32 s2, 0xffff
s_mov_b32 s3, 0x08027fac
v_ashrrev_i32 v1, 31, v0
s_load_dword s4, s[8:11], 0x04
v_lshrrev_b32 v1, 29, v1
v_add_i32 v1, vcc, v0, v1
v_and_b32 v1, -8, v1
v_lshlrev_b32 v2, 2, v0
// v_mov_b32 v2, 4
v_sub_i32 v1, vcc v0 v1
v_lshlrev_b32 v0, 2, v0
//v_cvt_f32_i32 v2, v2
v_lshlrev_b32 v1, 2, v1
ds_write_b32 v0, v0, v2
s_waitcnt lgkmcnt(0)
v_add_i32 v0, vcc, s4, v0
ds_read_b32 v1, v1
v_cmp_gt_u32 vcc, -1, v0
s_and_saveexec_b64 s[8:11], vcc
s_waitcnt lgkmcnt(0)
buffer_store_dword v1, v0, s[0:3], 0, offen
s_endpgm
};";
// Initialize OpenClWithGCN
OpenClWithGCN gprog = new OpenClWithGCN();
OpenClEnvironment env = gprog.env; // use the default environment
bool success = gprog.GcnCompile(source);
Console.Write(env.lastMessage);
if (!success)
{
return;
}
// Create some random data
var random = new Random();
const int count = 256;
const int dataSz = count * sizeof(float);
float[] data = (from i in Enumerable.Range(0, count)
select(float) random.NextDouble()).ToArray();
// Create a kernel from our modProgram
Kernel kernel = env.program.CreateKernel("myAsmFunc");
// Allocate an input and output memory buffers
Mem input = env.context.CreateBuffer(MemoryFlags.ReadOnly, dataSz);
Mem output = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);
// Copy our host buffer of random values to input device buffer
env.cmdQueue.EnqueueWriteBuffer(input, true, 0, dataSz, data);
// Set the arguments to our kernel, and enqueue it for execution
kernel.Arguments[0].SetValue(input);
kernel.Arguments[1].SetValue(output);
// Enqueue and run the kernel.
env.cmdQueue.EnqueueNDRangeKernel(kernel, count, 256);
// Wait until all commands finish.
env.cmdQueue.Finish();
// Read back the results.
int[] results = new int[count];
env.cmdQueue.EnqueueReadBufferAndWait(output, results);
// Print information
//for (int i = 0; i < 256; i++)
// Console.WriteLine("Debug: {0}: {1}", i, results[i]);
// Validate and print a brief summary detailing the results
int correct = 0;
for (int i = 0; i < count; i++)
{
correct += (results[i] == ((i * 4) % 32)) ? 1 : 0;
}
Console.WriteLine("{0} - Computed {1}/{2} correct values!",
correct == count ? "PASS" : "FAIL", correct, count);
Assert.AreEqual(correct, count);
}
public void Run()
{
Console.WriteLine("================== " + GetType().Name + "================== ");
string source = @"__asm4GCN myAsmFunc (float*,float*,float*)
{
// ======== Pre-Loaded Registers =========
// s[2:3] - PTR_UAV_TABLE
// s[2:3] +0x60 - base_resource_const1(#T)
// s[2:3] +0x68 - base_resource_const2(#T)
// s[4:7] - IMM_CONST_BUFFER0
// s[4:7] +0x00 - Grid Size
// s[4:7] +0x04 - Local Size
// s[4:7] +0x18 - baseGlobalId
// s[8:11] -IMM_CONST_BUFFER1
// s[8:11]+0x00 - param1 offset
// s[8:11]+0x04 - param2 offset
// s[8:11]+0x08 - param3 offset
// s1 - threadgroupId
// s12 - groupId
// v0 - laneId
// Use a #define to shorten tbuffer_load instruction
#define _F32_ 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT]
// Variable decelerations with specific registers (pre-loaded regs)
s8u resConsts s[2:3] // base resource constants
s16u uavBuff s[4:7], paramsOffset s[8:11], groupId s12
v4b laneId v0 // v0 is pre-loaded with lane id
// Setup Resource Constant
s_load_dwordx4 s16u resConst1, resConsts, 0x68
// establish GlobalID
s_buffer_load_dword s4u baseGlobalId,uavBuff, 0x18
s_buffer_load_dword s4u localSize, uavBuff, 0x04
s_waitcnt lgkmcnt(0)
v_mov_b32 v4u vLocalSize, localSize
v_mul_i32_i24 vLocalSize, groupId, vLocalSize
v_add_i32 v4u localSizeIdx, vcc, laneId, vLocalSize
v_add_i32 v4u vGlobalID, vcc, baseGlobalId, localSizeIdx
v_lshlrev_b32 v4u vGlobalOffset, 2, vGlobalID
// create a tool called 'ShowVar' to easily output a variable
#define _ShowVar_(VarToPrint) \
v_mov_b32 v4u debugToPrint, VarToPrint;\
tbuffer_store_format_x debugToPrint, debugOffset,resConst1,_F32_
s_buffer_load_dword s4u param3Offset, paramsOffset, 0x08
s_waitcnt lgkmcnt(0)
v_add_i32 v4u debugOffset, vcc, param3Offset, vGlobalOffset
// Fetch the value that we want to multiply by 2.
s_buffer_load_dword s4u sPara1Ptr, paramsOffset, 0x00
s_load_dwordx4 s16u resConst0, resConsts, 0x60
s_waitcnt lgkmcnt(0)
v_add_i32 v4b baseOffset, vcc, sPara1Ptr, vGlobalOffset
tbuffer_load_format_x v4b val, baseOffset, resConst0, _F32_
s_waitcnt vmcnt(0)
// Fast 64-lane wavefront SUM reduction
// Product, Avg, Min, Max can also be easily achieved with minor edits.
v4f tmp;
[[ for (int i = 2; i <7; i++) {]]
ds_swizzle_b32 tmp, val, tmp, tmp offset1:[[= 1<<i ]] offset0:0b00011111
s_waitcnt lgkmcnt(0)
v_add_f32 val, tmp, val // can also use v_min_f32, v_max_f32, or v_mul_f32
[[ } ]]
v_readfirstlane_b32 s4u sum, val
v_add_f32 val, sum, val
v_readlane_b32 sum, val, 32
// Write the results back to memory.
v_mov_b32 val, sum
s_buffer_load_dword s4u param2Offset, paramsOffset, 0x04
s_waitcnt lgkmcnt(0)
s_waitcnt vmcnt(0)
v_add_i32 v4u dstOffset, vcc, param2Offset,vGlobalOffset
free param2Offset // var is freed and register returned to pool.
tbuffer_store_format_x val, dstOffset,resConst1,_F32_
_ShowVar_(val)
// Exit the kernel
s_endpgm
};";
// https://visualstudiogallery.msdn.microsoft.com/46c0c49e-f825-454b-9f6a-48b216797eb5/view/Reviews/0?showReviewForm=True
// Initialize OpenClWithGCN
OpenClWithGCN gprog = new OpenClWithGCN();
OpenClEnvironment env = gprog.env; // use the default environment
bool success = gprog.GcnCompile(source);
Console.Write(env.lastMessage);
if (!success)
{
return;
}
// Create some random data
var random = new Random(3);
const int count = 640000;
const int dataSz = count * sizeof(float);
float[] data = (from i in Enumerable.Range(0, count)
select(float) random.NextDouble()).ToArray();
// Create a kernel from our modProgram
Kernel kernel = env.program.CreateKernel("myAsmFunc");
// Allocate an input and output memory buffers
Mem input = env.context.CreateBuffer(MemoryFlags.ReadOnly, dataSz);
Mem output = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);
Mem debug = env.context.CreateBuffer(MemoryFlags.WriteOnly, dataSz);
// Copy our host buffer of random values to input device buffer
env.cmdQueue.EnqueueWriteBuffer(input, true, 0, dataSz, data);
// Set the arguments to our kernel, and enqueue it for execution
kernel.Arguments[0].SetValue(input);
kernel.Arguments[1].SetValue(output);
kernel.Arguments[2].SetValue(debug);
// Enqueue and run the kernel.
env.cmdQueue.EnqueueNDRangeKernel(kernel, count, 256);
// Wait until all commands finish.
env.cmdQueue.Finish();
// Read back the results.
float[] results = new float[count];
env.cmdQueue.EnqueueReadBufferAndWait(output, results);
float[] debugOutput = new float[count];
env.cmdQueue.EnqueueReadBufferAndWait(debug, debugOutput);
// Print Debug information
if (printDetails)
{
for (int i = 0; i < 256; i++)
{
Console.WriteLine("Debug: {0}: {1} -> {2}", i, data[i], debugOutput[i]);
}
}
// Validate and print a brief summary detailing the results
int correct = 0;
for (int i = 0; i < (count / 64); i++)
{
float cpuSum = 0;
for (int j = 0; j < 64; j++)
{
cpuSum += data[i * 64 + j];
}
float gpuSum = results[i * 64];
if (cpuSum > gpuSum * 0.999999 &&
cpuSum < gpuSum * 1.000001)
{
correct++;
}
}
Console.WriteLine("Computed {0}/{1} correct values!", correct, count / 64);
Assert.AreEqual(correct, count / 64);
}
private void showSpecsInNotepadToolStripMenuItem_Click(object sender, EventArgs e)
{
NotepadHelper.ShowMessage(OpenClWithGCN.GetGPUInfo());
}