AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

ECR #333755 - Fixed a typo.

Browse Source 
[git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1132109]
This commit is contained in:
Ding, Wei (xN/A) TX
2015-03-18 17:32:10 -05:00
parent 85acd18466
commit 3a8bc1b01a
8 changed files with 0 additions and 542 deletions
Download Patch File Download Diff File Expand all files Collapse all files
samples/BitionicSort/Makefile
-7
View File
@@ -1,7 +0,0 @@
OPENCL_DEPTH = ../../..
include $(OPENCL_DEPTH)/runtimenew/runtimedefs
SUBDIRS = build
include $(OPENCL_DEPTH)/runtimenew/runtimerules
samples/BitionicSort/bitonic_sort.cpp
-215
View File
@@ -1,215 +0,0 @@
#include "bitonic_sort.h"
BitonicSort::BitonicSort()
{
}
BitonicSort::~BitonicSort()
{
}
void BitonicSort::InitlizeData()
{
sort_flag = 1;
length = 256;
verification_input = (uint *) malloc(length * sizeof(uint));
int input_size_bytes = length * sizeof(uint);
input_array = (uint*)malloc(input_size_bytes);
FillRandom(input_array, length, 1, 0, 255, 16);
memcpy(verification_input, input_array, length * sizeof(uint));
}
int BitonicSort::FillRandom(uint * arrayPtr, const int width, const int height, const uint rangeMin, const uint rangeMax, unsigned int seed)
{
if(!arrayPtr)
{
printf("Cannot fill array. NULL pointer.\n");
return -1;
}
if(!seed)
seed = (unsigned int)time(NULL);
srand(seed);
double range = double(rangeMax - rangeMin) + 1.0;
/* random initialisation of input */
for(int i = 0; i < height; i++)
for(int j = 0; j < width; j++)
{
int index = i*width + j;
arrayPtr[index] = rangeMin + (uint)(range*rand()/(RAND_MAX + 1.0));
}
return 0;
}
int BitonicSort::VerifyResults()
{
BitonicSortCPUReference(verification_input, length, sort_flag);
#ifdef DEBUG
for (int i=0; i<length; ++i)
{
printf("%d ", verification_input[i]);
}
printf("\n\n");
#endif
// compare the results and see if they match
if(memcmp(input_array, verification_input, length*sizeof(uint)) == 0)
{
cout<<"PASSED!\n" << endl;
return 0;
}
else
{
cout<<"FAILED\n" << endl;
return -1;
}
return 0;
}
void BitonicSort::BitonicSortCPUReference(uint * input, const uint length, const bool sortIncreasing)
{
const uint halfLength = length/2;
uint i;
for(i = 2; i <= length; i *= 2)
{
uint j;
for(j = i; j > 1; j /= 2)
{
bool increasing = sortIncreasing;
const uint half_j = j/2;
uint k;
for(k = 0; k < length; k += j)
{
const uint k_plus_half_j = k + half_j;
uint l;
if(i < length)
{
if((k == i) || ((k % i) == 0) && (k != halfLength))
increasing = !increasing;
}
for(l = k; l < k_plus_half_j; ++l)
{
if(increasing)
SwapIfFirstIsGreater(&input[l], &input[l + half_j]);
else
SwapIfFirstIsGreater(&input[l + half_j], &input[l]);
}
}
}
}
}
void BitonicSort::SwapIfFirstIsGreater(uint *a, uint *b)
{
if(*a > *b)
{
uint temp = *a;
*a = *b;
*b = temp;
}
}
void BitonicSort::RunKernels()
{
int num_of_stages = 0;
for(uint temp = length; temp > 1; temp >>= 1)
++num_of_stages;
for(int stage = 0; stage < num_of_stages; ++stage)
{
// Every stage has stage + 1 passes
for(int pass_of_stage = 0; pass_of_stage < stage + 1; ++pass_of_stage)
{
SetStages(stage, pass_of_stage);
}
}
}
void BitonicSort::Clean()
{
free(input_array);
input_array = NULL;
free(verification_input);
verification_input = NULL;
}
void BitonicSort::SetStages(uint num_of_stage, uint pass_of_stage)
{
struct __attribute__ ((aligned(HSA_ARGUMENT_ALIGN_BYTES))) args_t {
uint* offset_0;
uint* offset_1;
uint* offset_2;
uint* printf_buffer;
uint* vqueue_buffer;
uint* aqlwrap_pointer;
uint* inpu_array;
uint stage;
uint pass_of_stage;
uint direction;
} local_args;
local_args.offset_0 = 0;
local_args.offset_1 = 0;
local_args.offset_2 = 0;
local_args.printf_buffer = 0;
local_args.vqueue_buffer = 0;
local_args.aqlwrap_pointer = 0;
local_args.inpu_array = input_array;
local_args.stage = num_of_stage;
local_args.pass_of_stage = pass_of_stage;
local_args.direction = sort_flag;
///////////////////////
int group_x = GROUP_SIZE;
int group_y = 1;
int group_z = 1;
int group_size = 0;
int grid_x = length / 2 ;
int grid_y = 1;
int grid_z = 1;
int kernel_size = sizeof(args_t);
Run(1, group_x, group_y, group_z, group_size, grid_x, grid_y, grid_z, &local_args, kernel_size);
}
int main(int argc, char *argv[])
{
//char file_name[128] = "bitonic_sort_kernel.brig";
char file_name[128] = "bitonic_sort_kernel.hsail";
char kernel_name[128] = "&__OpenCL_bitonicSort_kernel";
BitonicSort bitonic;
//bitonic.SetBrigFileAndKernelName(file_name, kernel_name);
bitonic.GetHsailNameAndKernelName(file_name, kernel_name);
bitonic.InitlizeData();
bitonic.HsaInit();
bitonic.RunKernels();
bitonic.VerifyResults();
bitonic.Clean();
bitonic.Close();
return 0;
}
samples/BitionicSort/bitonic_sort.h
-54
View File
@@ -1,54 +0,0 @@
#ifndef __HSA_BITONIC_SORT__
#define __HSA_BITONIC_SORT__
#include "utilities.h"
#include <vector>
#include "hsa.h"
#include "hsa_ext_finalize.h"
#include "utilities.h"
#include <string.h>
#include<iostream>
using namespace std;
#include "hsa_base_test.h"
#define GROUP_SIZE 256
class BitonicSort : public HSA_TEST
{
public:
BitonicSort();
~BitonicSort();
public:
void SetStages(uint num_of_stage, uint pass_of_stage);
int FillRandom(uint * arrayPtr, const int width, const int height, const uint rangeMin, const uint rangeMax, unsigned int seed);
void RunKernels();
int VerifyResults();
void BitonicSortCPUReference(uint * input, const uint length, const bool sortIncreasing) ;
void SwapIfFirstIsGreater(uint *a, uint *b);
void InitlizeData();
void Clean();
public:
int length;
int sort_flag;
uint* verification_input;
uint* input_array;
uint* outBuffer_;
unsigned int width_;
unsigned int height_;
unsigned int bufSize_;
unsigned int blockSize_;
//static const unsigned int MAX_ITERATIONS = 50;
};
#endif
samples/BitionicSort/bitonic_sort_kernel.brig
BIN
View File
Binary file not shown.
samples/BitionicSort/bitonic_sort_kernel.cl
-45
View File
@@ -1,45 +0,0 @@
__kernel void bitonicSort(__global uint * theArray, const uint stage, const uint passOfStage, const uint direction)
{
uint sortIncreasing = direction;
uint threadId = get_global_id(0);
uint pairDistance = 1 << (stage - passOfStage);
uint blockWidth = 2 * pairDistance;
uint leftId = (threadId % pairDistance) + (threadId / pairDistance) * blockWidth;
uint rightId = leftId + pairDistance;
uint leftElement = theArray[leftId];
uint rightElement = theArray[rightId];
uint sameDirectionBlockWidth = 1 << stage;
if((threadId/sameDirectionBlockWidth) % 2 == 1)
sortIncreasing = 1 - sortIncreasing;
uint greater;
uint lesser;
if(leftElement > rightElement)
{
greater = leftElement;
lesser = rightElement;
}
else
{
greater = rightElement;
lesser = leftElement;
}
if(sortIncreasing)
{
theArray[leftId] = lesser;
theArray[rightId] = greater;
}
else
{
theArray[leftId] = greater;
theArray[rightId] = lesser;
}
}
samples/BitionicSort/bitonic_sort_kernel.hsail
-97
View File
@@ -1,97 +0,0 @@
version 0:20140528:$full:$large;
extension "amd:gcn";
extension "IMAGE";
decl prog function &abort()();
prog kernel &__OpenCL_bitonicSort_kernel(
kernarg_u64 %__global_offset_0,
kernarg_u64 %__global_offset_1,
kernarg_u64 %__global_offset_2,
kernarg_u64 %__printf_buffer,
kernarg_u64 %__vqueue_pointer,
kernarg_u64 %__aqlwrap_pointer,
kernarg_u64 %theArray,
kernarg_u32 %stage,
kernarg_u32 %passOfStage,
kernarg_u32 %direction)
{
pragma "AMD RTI", "ARGSTART:__OpenCL_bitonicSort_kernel";
pragma "AMD RTI", "version:3:1:104";
pragma "AMD RTI", "device:generic";
pragma "AMD RTI", "uniqueid:1024";
pragma "AMD RTI", "memory:private:0";
pragma "AMD RTI", "memory:region:0";
pragma "AMD RTI", "memory:local:0";
pragma "AMD RTI", "value:__global_offset_0:u64:1:1:0";
pragma "AMD RTI", "value:__global_offset_1:u64:1:1:16";
pragma "AMD RTI", "value:__global_offset_2:u64:1:1:32";
pragma "AMD RTI", "pointer:__printf_buffer:u8:1:1:48:uav:7:1:RW:0:0:0";
pragma "AMD RTI", "value:__vqueue_pointer:u64:1:1:64";
pragma "AMD RTI", "value:__aqlwrap_pointer:u64:1:1:80";
pragma "AMD RTI", "pointer:theArray:u32:1:1:96:uav:7:4:RW:0:0:0";
pragma "AMD RTI", "value:stage:u32:1:1:112";
pragma "AMD RTI", "constarg:7:stage";
pragma "AMD RTI", "value:passOfStage:u32:1:1:128";
pragma "AMD RTI", "constarg:8:passOfStage";
pragma "AMD RTI", "value:direction:u32:1:1:144";
pragma "AMD RTI", "constarg:9:direction";
pragma "AMD RTI", "function:1:0";
pragma "AMD RTI", "memory:64bitABI";
pragma "AMD RTI", "privateid:8";
pragma "AMD RTI", "enqueue_kernel:0";
pragma "AMD RTI", "kernel_index:0";
pragma "AMD RTI", "reflection:0:size_t";
pragma "AMD RTI", "reflection:1:size_t";
pragma "AMD RTI", "reflection:2:size_t";
pragma "AMD RTI", "reflection:3:size_t";
pragma "AMD RTI", "reflection:4:size_t";
pragma "AMD RTI", "reflection:5:size_t";
pragma "AMD RTI", "reflection:6:uint*";
pragma "AMD RTI", "reflection:7:uint";
pragma "AMD RTI", "reflection:8:uint";
pragma "AMD RTI", "reflection:9:uint";
pragma "AMD RTI", "ARGEND:__OpenCL_bitonicSort_kernel";
@__OpenCL_bitonicSort_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%__global_offset_0];
add_u64 $d0, $d0, $d1;
cvt_u32_u64 $s1, $d0;
ld_kernarg_align(4)_width(all)_u32 $s2, [%stage];
shl_u32 $s0, 1, $s2;
and_b32 $s0, $s1, $s0;
ld_kernarg_align(4)_width(all)_u32 $s3, [%direction];
sub_u32 $s4, 1, $s3;
cmp_eq_b1_s32 $c0, $s0, 0;
cmov_b32 $s0, $c0, $s3, $s4;
ld_kernarg_align(4)_width(all)_u32 $s3, [%passOfStage];
sub_u32 $s4, $s2, $s3;
shl_u32 $s2, 1, $s4;
add_u32 $s3, $s2, 4294967295;
and_b32 $s3, $s1, $s3;
shr_u32 $s1, $s1, $s4;
mul_u32 $s1, $s2, $s1;
shl_u32 $s1, $s1, 1;
ld_kernarg_align(8)_width(all)_u64 $d1, [%theArray];
add_u32 $s3, $s1, $s3;
cvt_u64_u32 $d0, $s3;
shl_u64 $d0, $d0, 2;
add_u64 $d0, $d1, $d0;
ld_global_align(4)_u32 $s1, [$d0];
add_u32 $s2, $s3, $s2;
cvt_u64_u32 $d2, $s2;
shl_u64 $d2, $d2, 2;
add_u64 $d1, $d1, $d2;
ld_global_align(4)_u32 $s3, [$d1];
max_u32 $s2, $s1, $s3;
min_u32 $s1, $s1, $s3;
cmp_eq_b1_s32 $c0, $s0, 0;
cmov_b32 $s0, $c0, $s1, $s2;
cmov_b32 $s1, $c0, $s2, $s1;
st_global_align(4)_u32 $s1, [$d0];
st_global_align(4)_u32 $s0, [$d1];
ret;
};
samples/BitionicSort/build/Makefile
-8
View File
@@ -1,8 +0,0 @@
OPENCL_DEPTH = ../../../..
include $(OPENCL_DEPTH)/runtimenew/runtimedefs
BUILD_SUBDIRS = $(DEFAULT_TARGETS)
BUILD_MAKEFILE = Makefile.BitionicSort
include $(OPENCL_DEPTH)/runtimenew/runtimerules
samples/BitionicSort/build/Makefile.BitionicSort
-116
View File
@@ -1,116 +0,0 @@
include $(OPENCL_DEPTH)/hsadefs
# Executable containing all the API core tests
EXE_TARGET = BitionicSort
#
# Makefile construct to copy executable to Dist folder
#
DIST_BIN_TARGET = $(EXE_TARGET)
ifdef ATI_OS_WINDOWS
DIST_BIN_TARGETS = $(EXE_TARGET)$(EXE_EXT)
endif
ifdef ATI_BITS_64
LIB_SUFFIX=64
BITS=64
CMPILERBITS=64
else
LIB_SUFFIX=
BITS=
ifdef ATI_OS_WINDOWS
CMPILERBITS=
else
CMPILERBITS=32
endif
endif
ifdef ATI_OS_WINDOWS
LFLAGS += /subsystem:console
CORE_LIB=dll
LIB_PREFIX=
DYN_LIB_EXT=$(LIB_EXT)
UTIL_EXT=win
else
LIB_PREFIX=lib
CORE_LIB=so
DYN_LIB_EXT=$(DLL_EXT)
UTIL_EXT=
endif
LCXXDEFS += -DAMD_INTERNAL_BUILD
vpath %.cpp $(COMPONENT_DEPTH)
CPPFILES := $(notdir $(wildcard $(COMPONENT_DEPTH)/*.cpp))
TOOLS_TEST_COMMON=$(COMPONENT_DEPTH)/../common
vpath %.cpp $(TOOLS_TEST_COMMON)
CPPFILES += $(notdir $(wildcard $(TOOLS_TEST_COMMON)/*.cpp))
ifdef ATI_OS_LINUX
GCXXOPTS := $(filter-out -fno-rtti,$(GCXXOPTS))
GCXXOPTS := $(filter-out -fno-exceptions,$(GCXXOPTS))
LFLAGS += -L$(DIST_LIB_DEST) -lpthread $(LIBSTDCXX) -lm -ldl -lrt -lstdc++
#LFLAGS += -L$(DIST_LIB_DEST) -lpthread $(LIBSTDCXX) -lm -ldl -lrt
endif
#LCINCS := $(INCSWITCH) "$(OPENCL_DEPTH)/contrib/gtest-1.6.0/include"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/runtime"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/runtime/samples"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/contrib/elftoolchain/libelf"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/contrib/elftoolchain/common"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/runtime/inc"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/Interface"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL/$(FULL_BUILD_DIR)"
LCINCS += $(INCSWITCH) "$(OPENCL_PATH)/compiler/lib/include"
LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/runtime/samples/common"
#HSAIL Assembler header files
#LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL"
#LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL/$(FULL_BUILD_DIR)"
#LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/runtime/tools/sp3/Chip/R1000"
#LLLIBS := $(OPENCL_DEPTH)/contrib/gtest-1.6.0/$(FULL_BUILD_DIR)/libgtest$(LIB_EXT)
#LLLIBS += $(OPENCL_DEPTH)/runtime/test/common/$(FULL_BUILD_DIR)/testcommon$(LIB_EXT)
#LLLIBS += $(OPENCL_DEPTH)/runtime/test/gcommon/$(FULL_BUILD_DIR)/gtestcommon$(LIB_EXT)
#LLLIBS += $(OPENCL_DEPTH)/runtime/tools/sp3/Chip/R1000/$(FULL_BUILD_DIR)/SP3_R1000$(LIB_EXT)
#LLLIBS += $(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL/$(FULL_BUILD_DIR)/libhsail$(LIB_EXT)
UTIL_BUILD=build/$(OS_TYPE)/util/$(UTIL_EXT)/$(BUILD_DIR)
#LLLIBS += $(OPENCL_DEPTH)/runtime/core/$(UTIL_BUILD)/util$(UTIL_EXT)$(LIB_EXT)
LIBELF_DIR=build/$(OS_TYPE)/$(BUILD_DIR)
RUNTIME_BUILD=build/$(OS_TYPE)/$(CORE_LIB)/$(BUILD_DIR)
ifdef ATI_OS_LINUX
GCXXOPTS := $(filter-out -fno-exceptions,$(GCXXOPTS))
LFLAGS += -L$(DIST_LIB_DEST) -lpthread $(LIBSTDCXX) -lm -ldl -lrt
LFLAGS += -L$(OPENCL_DEPTH)/contrib/elftoolchain/libelf/$(LIBELF_DIR) -lelf
LFLAGS += -L$(OPENCL_DEPTH)/runtime/core/$(RUNTIME_BUILD) -lhsa-runtime$(LIB_SUFFIX)
LFLAGS += -L$(OPENCL_DEPTH)/runtime/tools/$(RUNTIME_BUILD) -lhsa-runtime-tools$(LIB_SUFFIX)
LLLIBS = $(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL/$(FULL_BUILD_DIR)/libhsail$(LIB_EXT)
else
# Verify the extension of libelf is valid i.e. is not dll but instead is "lib"
#LLLIBS += -L$(OPENCL_DEPTH)/contrib/elftoolchain/libelf/$(LIBELF_DIR)/libelf
LLLIBS += $(OPENCL_DEPTH)/runtime/core/$(RUNTIME_BUILD)/hsa-runtime$(LIB_SUFFIX)$(LIB_EXT)
LLLIBS += $(OPENCL_DEPTH)/runtime/tools/$(RUNTIME_BUILD)/hsa-runtime-tools$(LIB_SUFFIX)$(LIB_EXT)
endif
#
# Copy the Hsa ample to Hsa distribution folder
#
ifdef ATI_OS_LINUX
LOCAL_TARGETS = inst_sample_copy
endif
inst_sample_copy:
$(MKDIR) $(DIST_BIN_DEST)/../../samples/$(OS_TYPE)
-$(RM) $(DIST_BIN_DEST)/../../samples/$(OS_TYPE)/bitonic_sort_kernel.hsail
$(CP) $(BUILD_DIR)/$(EXE_TARGET) $(DIST_BIN_DEST)/../../samples/$(OS_TYPE)
$(CP) $(COMPONENT_DEPTH)/bitonic_sort_kernel.hsail $(DIST_BIN_DEST)/../../samples/$(OS_TYPE)
include $(OPENCL_DEPTH)/hsarules