diff --git a/samples/BitionicSort/Makefile b/samples/BitionicSort/Makefile deleted file mode 100644 index 11141a08a8..0000000000 --- a/samples/BitionicSort/Makefile +++ /dev/null @@ -1,7 +0,0 @@ -OPENCL_DEPTH = ../../.. - -include $(OPENCL_DEPTH)/runtimenew/runtimedefs - -SUBDIRS = build - -include $(OPENCL_DEPTH)/runtimenew/runtimerules diff --git a/samples/BitionicSort/bitonic_sort.cpp b/samples/BitionicSort/bitonic_sort.cpp deleted file mode 100644 index e30451efc7..0000000000 --- a/samples/BitionicSort/bitonic_sort.cpp +++ /dev/null @@ -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 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; -} - - diff --git a/samples/BitionicSort/bitonic_sort.h b/samples/BitionicSort/bitonic_sort.h deleted file mode 100644 index cc1ac8bb0a..0000000000 --- a/samples/BitionicSort/bitonic_sort.h +++ /dev/null @@ -1,54 +0,0 @@ -#ifndef __HSA_BITONIC_SORT__ -#define __HSA_BITONIC_SORT__ - - -#include "utilities.h" -#include - -#include "hsa.h" -#include "hsa_ext_finalize.h" -#include "utilities.h" - -#include -#include -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 - diff --git a/samples/BitionicSort/bitonic_sort_kernel.brig b/samples/BitionicSort/bitonic_sort_kernel.brig deleted file mode 100644 index 4ed953e993..0000000000 Binary files a/samples/BitionicSort/bitonic_sort_kernel.brig and /dev/null differ diff --git a/samples/BitionicSort/bitonic_sort_kernel.cl b/samples/BitionicSort/bitonic_sort_kernel.cl deleted file mode 100644 index 94cd8cc9f6..0000000000 --- a/samples/BitionicSort/bitonic_sort_kernel.cl +++ /dev/null @@ -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; - } -} - diff --git a/samples/BitionicSort/bitonic_sort_kernel.hsail b/samples/BitionicSort/bitonic_sort_kernel.hsail deleted file mode 100644 index 893a389894..0000000000 --- a/samples/BitionicSort/bitonic_sort_kernel.hsail +++ /dev/null @@ -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; -}; diff --git a/samples/BitionicSort/build/Makefile b/samples/BitionicSort/build/Makefile deleted file mode 100644 index 2fc287addd..0000000000 --- a/samples/BitionicSort/build/Makefile +++ /dev/null @@ -1,8 +0,0 @@ -OPENCL_DEPTH = ../../../.. - -include $(OPENCL_DEPTH)/runtimenew/runtimedefs - -BUILD_SUBDIRS = $(DEFAULT_TARGETS) -BUILD_MAKEFILE = Makefile.BitionicSort - -include $(OPENCL_DEPTH)/runtimenew/runtimerules diff --git a/samples/BitionicSort/build/Makefile.BitionicSort b/samples/BitionicSort/build/Makefile.BitionicSort deleted file mode 100644 index 5756a6c96f..0000000000 --- a/samples/BitionicSort/build/Makefile.BitionicSort +++ /dev/null @@ -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 -