minor update

.travis.yml

@@ -22,31 +22,35 @@ cache:
     - $HOME/build64
 
 before_install:
-  - if [ ! -d "$TOOLDIR" ] || [ -z "$(ls -A $TOOLDIR)" ]; then
-      rm -rf $TOOLDIR;
-      mkdir -p $TOOLDIR;
-      OSDIR=ubuntu/focal ./ci/toolchain_install.sh --all;
-    fi
-  - if [ ! -d "$HOME/third_party" ] || [ -z "$(ls -A $HOME/third_party)" ]; then
-      make -C third_party > /dev/null;
-      cp -r third_party $HOME/third_party;
-    else
-      rsync -v --ignore-existing $HOME/third_party/ third_party/;
-    fi
+  - cp -r $HOME/third_party third_party;
   - source ./ci/toolchain_env.sh
 
 stages:
   - setup
+  - build
   - test
 
 jobs:     
   include:
     - stage: setup
+      script:
+        - if [ ! -d "$TOOLDIR" ] || [ -z "$(ls -A $TOOLDIR)" ] || [ "$(cat "$TOOLDIR/version.txt")" != "v1" ]; then
+          rm -rf $TOOLDIR;
+          mkdir -p $TOOLDIR;
+          OSDIR=ubuntu/focal ./ci/toolchain_install.sh --all;
+          echo "v1" > "$TOOLDIR/version.txt";
+        fi
+      - if [ ! -d "$HOME/third_party" ] || [ -z "$(ls -A $HOME/third_party)" ] || [ "$(cat "$HOME/third_party/version.txt")" != "v1" ]; then
+          make -C third_party > /dev/null;
+          cp -r third_party $HOME/third_party;
+          echo "v1" > "$HOME/third_party/version.txt";          
+        fi
+    - stage: build
       name: build32
       script:
         - rm -rf $HOME/build32 && mkdir -p $HOME/build32
         - cd $HOME/build32 && $TRAVIS_BUILD_DIR/configure && make -s > /dev/null
-    - stage: setup
+    - stage: build
       name: build64
       script:
         - rm -rf $HOME/build64 && mkdir -p $HOME/build64

tests/opencl/common.mk

@@ -53,6 +53,9 @@ OBJS := $(addsuffix .o, $(filter-out main.cc,$(notdir $(SRCS))))
 
 all: $(PROJECT) kernel.pocl
  
+kernel.cl: $(SRC_DIR)/kernel.cl
+	cp $(SRC_DIR)/kernel.cl $@
+
 kernel.pocl: $(SRC_DIR)/kernel.cl
 	LD_LIBRARY_PATH=$(LLVM_POCL)/lib:$(POCL_CC_PATH)/lib:$(LLVM_VORTEX)/lib:$(LD_LIBRARY_PATH) LLVM_PREFIX=$(LLVM_VORTEX) POCL_DEBUG=all POCL_VORTEX_CFLAGS="$(K_CFLAGS)" POCL_VORTEX_LDFLAGS="$(K_LDFLAGS)" $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl $^
  
@@ -78,10 +81,10 @@ endif
 $(PROJECT): setup main.cc.o $(OBJS)
 	$(CXX) $(CXXFLAGS) $(filter-out setup, $^) $(LDFLAGS) -L$(ROOT_DIR)/runtime/stub -lvortex -L$(POCL_RT_PATH)/lib -lOpenCL -o $@
 
-$(PROJECT).host: setup main.cc.host.o $(OBJS)
+$(PROJECT).host: setup main.cc.host.o $(OBJS)	
 	$(CXX) $(CXXFLAGS) $(filter-out setup, $^) $(LDFLAGS) -lOpenCL -o $@
 
-run-gpu: $(PROJECT).host kernel.pocl
+run-gpu: $(PROJECT).host kernel.cl
 	./$(PROJECT).host $(OPTS)
 
 run-simx: $(PROJECT) kernel.pocl   

tests/opencl/lbm/Makefile

@@ -8,7 +8,15 @@ SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
 SRCS := $(SRC_DIR)/main.cc $(SRC_DIR)/args.c $(SRC_DIR)/parboil_opencl.c $(SRC_DIR)/gpu_info.c $(SRC_DIR)/lbm.c $(SRC_DIR)/ocl.c
 
 CXXFLAGS += -I$(SRC_DIR)
-K_CFLAGS += -I$(SRC_DIR)
+
+lbm_macros.h: $(SRC_DIR)/lbm_macros.h
+	cp $< $@
+
+layout_config.h: $(SRC_DIR)/layout_config.h
+	cp $< $@
+
+setup: lbm_macros.h layout_config.h
+USE_SETUP := yes
 
 # Usage: #iter [-i input_file] [-o output_file]
 OPTS ?= 1 -i $(SRC_DIR)/32_32_8_ldc.of

tests/opencl/sgemm/Makefile

@@ -8,7 +8,12 @@ SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
 SRCS := $(SRC_DIR)/main.cc
 
 CXXFLAGS += -I$(SRC_DIR)
-K_CFLAGS += -I$(SRC_DIR)
+
+common.h: $(SRC_DIR)/common.h 
+	cp $< $@
+
+setup: common.h
+USE_SETUP := yes
 
 OPTS ?= -n32
 

tests/opencl/sgemm2/Makefile

@@ -7,6 +7,14 @@ SRC_DIR := $(VORTEX_HOME)/tests/opencl/$(PROJECT)
 
 SRCS := $(SRC_DIR)/main.cc
 
+CXXFLAGS += -I$(SRC_DIR)
+
+common.h: $(SRC_DIR)/common.h 
+	cp $< $@
+
+setup: common.h
+USE_SETUP := yes
+
 OPTS ?= -n32
 
 include ../common.mk

tests/opencl/sgemm2/kernel.cl

@@ -31,6 +31,9 @@ __kernel void sgemm2(__global float *A,
         for (int j = 0; j < LOCAL_SIZE; j++) {
             sum += localA[localRow][j] * localB[j][localCol];
         }
+
+        // Ensure computation is done before loading next block
+        barrier(CLK_LOCAL_MEM_FENCE);
     }
 
     C[globalRow * N + globalCol] = sum;