Skip to content

Instantly share code, notes, and snippets.

View hughperkins's full-sized avatar

Hugh Perkins hughperkins

1k followers · 7 following
View GitHub Profile
@hughperkins
hughperkins / cocl-issue3-c.cpp
Created November 5, 2016 14:23
#include "EasyCL/EasyCL.h"
#include <iostream>
#include <stdexcept>
#include "cocl/cocl.h"
using namespace std;
string kernelSource = R"(
@hughperkins
hughperkins / cocl-issue3-a.cpp
Created November 5, 2016 11:35
#include "EasyCL/EasyCL.h"
#include <iostream>
#include <stdexcept>
using namespace std;
string kernelSource = R"(
kernel void _z8setValuePfif(global float* data, long data_offset, int idx, float value) {
data = (global float*)((global char *)data + data_offset);
@hughperkins
hughperkins / gist:6d9d010225f8c9a5e28a7526c7063e78
Created November 5, 2016 11:21
import numpy as np
import pyopencl as cl
N = 32
its = 3
np.random.seed(123)
a = np.random.rand(N).astype(np.float32)
gpu_idx = 0
@hughperkins
hughperkins / cocl-issue3-a.cpp
Created November 5, 2016 09:00
#include "CL/cl.h"
#include <iostream>
#include <stdexcept>
using namespace std;
string kernelSource = R"(
kernel void _z8setValuePfif(global float* data, long data_offset, int idx, float value) {
data = (global float*)((global char *)data + data_offset);
@hughperkins
hughperkins / gist:dfe091c2637662c1ef75345c33c64ae6
Created November 3, 2016 23:40
import numpy as np
import pyopencl as cl
N = 32
its = 3
a = np.random.rand(N).astype(np.float32)
gpu_idx = 0
@hughperkins
hughperkins / gist:56225890541ff5404e4bcd69f93372a2
Created November 2, 2016 12:12
kernel void _ZN5Eigen8internal15EigenMetaKernelINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi1EiEELi16ENS_11MakePointerEEEKNS_18TensorCwiseUnaryOpINS0_14scalar_sqrt_opIfEEKNS4_INS5_IKfLi1ELi1EiEELi16ES7_EEEEEENS_9GpuDeviceEEEiEEvT_T0_(
global char* eval_nopointers, global float* eval_ptr0, long eval_ptr0_offset, global float* eval_ptr1, long eval_ptr1_offset, int size, local int *scratch) {
int i1_02_i;
int i_03_i;
int v2;
int v5;
int v7;
int v9;
int v11;
int v23;
@hughperkins
hughperkins / gist:d28e970d8df7e270a6fd68d67863318f
Created October 29, 2016 10:13
$ objdump -r build/test-cocl-cuda_sample.o | grep R_
000000000000000f R_X86_64_32 s.build/test-cocl-cuda_sample-device.cl._Z8setValuePfif
0000000000000014 R_X86_64_32 __opencl_sourcecodebuild/test-cocl-cuda_sample-device.cl
0000000000000019 R_X86_64_PC32 configureKernel-0x0000000000000004
0000000000000021 R_X86_64_PC32 _Z20setKernelArgCharStarPc-0x0000000000000004
0000000000000028 R_X86_64_PC32 _Z17setKernelArgInt32i-0x0000000000000004
0000000000000033 R_X86_64_PC32 _Z17setKernelArgFloatf-0x0000000000000004
000000000000003e R_X86_64_PC32 _Z8kernelGov-0x0000000000000004
000000000000006e R_X86_64_PC32 cudaMalloc-0x0000000000000004
0000000000000089 R_X86_64_PC32 cudaConfigureCall-0x0000000000000004
@hughperkins
hughperkins / placement log
Created October 28, 2016 14:42
W: /job:localhost/replica:0/task:0/gpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] W: /job:localhost/replica
:0/task:0/gpu:0
Assign: /job:localhost/replica:0/task:0/gpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] Assign: /job:localhost/re
plica:0/task:0/gpu:0
W/read: /job:localhost/replica:0/task:0/gpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] W/read: /job:localhost/re
plica:0/task:0/gpu:0
W/Assign: /job:localhost/replica:0/task:0/gpu:0
@hughperkins
hughperkins / gist:cc427f7be96639577c53889168ebb1c4
Created October 28, 2016 07:30
inline float __shfl_down_3(float v0, int v1, int v2) {
local float mem[1024];
int tid = get_local_id(0);
int warpid = tid % 32;
int warpstart = tid - warpid;
mem[tid] = v0;
//barrier(CLK_LOCAL_MEM_FENCE);
int warpsrc = warpid + v1;
warpsrc = warpsrc >= 32 ? warpid : warpsrc;
@hughperkins
hughperkins / reduction sum 3 elements
Created October 28, 2016 07:13
inline float __shfl_down_3(float v0, int v1, int v2) {
local float mem[1024];
int tid = get_local_id(0);
int warpid = tid % 32;
int warpstart = tid - warpid;
mem[tid] = v0;
//barrier(CLK_LOCAL_MEM_FENCE);
int warpsrc = warpid + v1;
warpsrc = warpsrc >= 32 ? warpid : warpsrc;