Skip to content

Commit

Permalink
Fixed several issues with macOS
Browse files Browse the repository at this point in the history
  • Loading branch information
ProjectPhysX committed Apr 30, 2023
1 parent 0f7e621 commit c5406ef
Show file tree
Hide file tree
Showing 2 changed files with 3 additions and 14 deletions.
6 changes: 3 additions & 3 deletions make.sh
Original file line number Diff line number Diff line change
Expand Up @@ -3,8 +3,8 @@
mkdir -p bin # create directory for executable
rm -f ./bin/OpenCL-Benchmark # prevent execution of old version if compiling fails

g++ ./src/*.cpp -o ./bin/OpenCL-Benchmark -pthread -I./src/OpenCL/include -L./src/OpenCL/lib -lOpenCL # compile on Linux
#g++ ./src/*.cpp -o ./bin/OpenCL-Benchmark -pthread -I./src/OpenCL/include -framework OpenCL # compile on macOS
#g++ ./src/*.cpp -o ./bin/OpenCL-Benchmark -pthread -I./src/OpenCL/include -L/system/vendor/lib64 -lOpenCL # compile on Android
g++ ./src/*.cpp -o ./bin/OpenCL-Benchmark -std=c++17 -pthread -I./src/OpenCL/include -L./src/OpenCL/lib -lOpenCL # compile on Linux
#g++ ./src/*.cpp -o ./bin/OpenCL-Benchmark -std=c++17 -pthread -I./src/OpenCL/include -framework OpenCL # compile on macOS
#g++ ./src/*.cpp -o ./bin/OpenCL-Benchmark -std=c++17 -pthread -I./src/OpenCL/include -L/system/vendor/lib64 -lOpenCL # compile on Android

./bin/OpenCL-Benchmark "$@" # run OpenCL-Benchmark
11 changes: 0 additions & 11 deletions src/kernel.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,6 @@ string opencl_c_container() { return R( // ########################## begin of O
kernel void kernel_double(global float* data) {
double x = (double)get_global_id(0);
double y = (double)get_local_id(0);
#pragma unroll
for(uint i=0u; i<128u; i++) {
x = fma(y, x, y);
y = fma(x, y, x);
Expand All @@ -19,7 +18,6 @@ kernel void kernel_double(global float* data) {
kernel void kernel_float(global float* data) {
float x = (float)get_global_id(0);
float y = (float)get_local_id(0);
#pragma unroll
for(uint i=0u; i<512u; i++) {
x = fma(y, x, y);
y = fma(x, y, x);
Expand All @@ -31,7 +29,6 @@ kernel void kernel_float(global float* data) {
kernel void kernel_half(global float* data) {
half2 x = (half2)((float)get_global_id(0), (float)get_local_id(0));
half2 y = (half2)((float)get_local_id(0), (float)get_global_id(0));
#pragma unroll
for(uint i=0u; i<512u; i++) {
x = fma(y, x, y);
y = fma(x, y, x);
Expand All @@ -43,7 +40,6 @@ kernel void kernel_half(global float* data) {
kernel void kernel_long(global float* data) {
long x = (long)get_global_id(0);
long y = (long)get_local_id(0);
#pragma unroll
for(uint i=0u; i<8u; i++) {
x = (y*x)+y;
y = (x*y)+x;
Expand All @@ -54,7 +50,6 @@ kernel void kernel_long(global float* data) {
kernel void kernel_int(global float* data) {
int x = get_global_id(0);
int y = get_local_id(0);
#pragma unroll
for(uint i=0u; i<512u; i++) {
x = (y*x)+y;
y = (x*y)+x;
Expand All @@ -65,7 +60,6 @@ kernel void kernel_int(global float* data) {
kernel void kernel_short(global float* data) {
short2 x = as_short2((int)get_global_id(0));
short2 y = as_short2((int)get_local_id(0));
#pragma unroll
for(uint i=0u; i<128u; i++) {
x = (y*x)+y;
y = (x*y)+x;
Expand All @@ -76,7 +70,6 @@ kernel void kernel_short(global float* data) {
kernel void kernel_char(global float* data) {
char4 x = as_char4((int)get_global_id(0));
char4 y = as_char4((int)get_local_id(0));
#pragma unroll
for(uint i=0u; i<64u; i++) {
x = (y*x)+y;
y = (x*y)+x;
Expand All @@ -88,25 +81,21 @@ kernel void kernel_char(global float* data) {

kernel void kernel_coalesced_write(global float* data) {
const uint n = get_global_id(0);
#pragma unroll
for(uint i=0u; i<def_M; i++) data[i*def_N+n] = 0.0f; // coalesced write
}
kernel void kernel_coalesced_read(global float* data) {
const uint n = get_global_id(0);
float x = 0.0f;
#pragma unroll
for(uint i=0u; i<def_M; i++) x += data[i*def_N+n]; // coalesced read
data[n] = x;
}
kernel void kernel_misaligned_write(global float* data) {
const uint n = get_global_id(0);
#pragma unroll
for(uint i=0u; i<def_M; i++) data[n*def_M+i] = 0.0f; // misaligned write
}
kernel void kernel_misaligned_read(global float* data) {
const uint n = get_global_id(0);
float x = 0.0f;
#pragma unroll
for(uint i=0u; i<def_M; i++) x += data[n*def_M+i]; // misaligned read
data[n] = x;
}
Expand Down

0 comments on commit c5406ef

Please sign in to comment.