Skip to content

Instantly share code, notes, and snippets.

@eholk
eholk / gist:4230942
Created December 7, 2012 05:09
Disassembled inner loop from ATLAS sdot
=> 0x00007fffefc226b0 <+0>: movslq %edi,%rax
0x00007fffefc226b3 <+3>: vxorps %xmm0,%xmm0,%xmm0
0x00007fffefc226b7 <+7>: and $0xfffffffc,%edi
0x00007fffefc226ba <+10>: lea (%rsi,%rax,4),%r8
0x00007fffefc226be <+14>: movslq %edi,%rdi
0x00007fffefc226c1 <+17>: lea (%rsi,%rdi,4),%rdi
0x00007fffefc226c5 <+21>: cmp %rdi,%rsi
0x00007fffefc226c8 <+24>: je 0x7fffefc22754 <ATL_sdot_xp1yp1aXbX+164>
0x00007fffefc226ce <+30>: vxorps %xmm1,%xmm1,%xmm1
0x00007fffefc226d2 <+34>: mov %rcx,%rdx
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_13, texmode_independent
.address_size 32
// .globl _ZN9add_float16_cb9e1b436595b333_00E
@eholk
eholk / gist:4217726
Created December 5, 2012 17:31
LLVM/PTX Error from unsupported LLVM code
unexpected address space
UNREACHABLE executed at /usr/local/src/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1257!
0 llc 0x000000010a169a72 void* llvm::object_creator<std::vector<std::string, std::allocator<std::string> > >() + 10722
1 llc 0x000000010a16a029 void* llvm::object_creator<std::vector<std::string, std::allocator<std::string> > >() + 12185
2 libsystem_c.dylib 0x00007fff89c728ea _sigtramp + 26
3 libsystem_c.dylib 0x00007fbcd1c05130 _sigtramp + 18446743787154253920
4 llc 0x000000010a169cc6 void* llvm::object_creator<std::vector<std::string, std::allocator<std::string> > >() + 11318
5 llc 0x000000010a158eb2 void* llvm::object_creator<std::vector<char const*, std::allocator<char const*> > >() + 49634
6 llc 0x0000000109819436 llvm::Pass* llvm::callDefaultCtor<llvm::NVPTXAllocaHoisting>() + 161766
7 llc 0x0000000109818487 llvm::Pass* llvm::callDefaultCtor<llvm::NVPTXAllocaHoisting>() + 157751
@eholk
eholk / trivial-kernel.ll
Created December 5, 2012 17:30
LLVM code generated by Rust
; ModuleID = 'trivial-kernel.rc'
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64"
target triple = "x86_64-apple-darwin"
%shapes = type opaque
%"~enum intrinsic::TyDesc[#0]" = type { [16 x i8] }
%tydesc = type { i64, i64, void (i1*, i1*, %tydesc**, i8*)*, void (i1*, i1*, %tydesc**, i8*)*, void (i1*, i1*, %tydesc**, i8*)*, void (i1*, i1*, %tydesc**, i8*)*, i8*, i8* }
@_rust_crate_map_trivial-kernel_0.0_68a2c114141ca = global { i32, i8*, i64, [2 x i64] } { i32 1, i8* bitcast (void (i1*, { i64, %tydesc*, i8*, i8*, i8 } addrspace(1)*)* @_ZN7cleanup10annihilate17_bde425c43da4f4993_05E to i8*), i64 ptrtoint ([1 x { i64, i64 }]* @_rust_mod_map to i64), [2 x i64] [i64 ptrtoint (i64* @_rust_crate_map_core_0.5_c3ca5d77d81b46c1 to i64), i64 0] }
@shapes = external global %shapes
; ModuleID = 'trivial-kernel.bc'
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64"
target triple = "x86_64-apple-darwin"
%shapes = type opaque
%"~enum intrinsic::TyDesc[#0]" = type { [16 x i8] }
%tydesc = type { i64, i64, void (i1*, i1*, %tydesc**, i8*)*, void (i1*, i1*, %tydesc**, i8*)*, void (i1*, i1*, %tydesc**, i8*)*, void (i1*, i1*, %tydesc**, i8*)*, i8*, i8* }
;@_rust_crate_map_trivial-kernel_0.0_68a2c114141ca = global { i32, i8*, i64, [2 x i64] } { i32 1, i8* bitcast (void (i1*, { i64, %tydesc*, i8*, i8*, i8 } addrspace(1)*)* @_ZN7cleanup10annihilate17_bde425c43da4f4993_05E to i8*), i64 ptrtoint ([1 x { i64, i64 }]* @_rust_mod_map to i64), [2 x i64] [i64 ptrtoint (i64* @_rust_crate_map_core_0.5_c3ca5d77d81b46c1 to i64), i64 0] }
;@shapes = external global %shapes
@eholk
eholk / generate.rs
Created November 22, 2012 01:40
The Beginnings of a Matrix Library
use num::Num;
fn identity<T: Copy Num, M: matrix::BasicMatrix<T>>(N: uint) -> M {
matrix::create(N, N, |i, j| {
if i == j {
num::from_int::<T>(1)
}
else {
num::from_int(0)
}
;; original program
(module
(define (main)
(let ((x (vector-r r1 1 2 3 4))
(y (vector-r r2 1 2 3 4)))
(return 0))))
;; before infer-regions
#include "harlan.hpp"
cl::program g_prog(g_ctx.createAndBuildProgramFromSource("#include \"rt/gpu_common.h\"\n\n#include \"rt/gpu_only.h\"\n\n__kernel void kernel_84(region_ptr kern_66, region_ptr A_1, void __global * mainregion_73_85) {\n region __global * mainregion_73 = ((region __global *)(mainregion_73_85));\n {\n int __global * retval_70 = (&(((int __global *)(get_region_ptr(mainregion_73, (kern_66) + (sizeof(int)))))[get_global_id(0)]));\n int __global * retval_69 = (&(((int __global *)(get_region_ptr(mainregion_73, (*retval_70) + (sizeof(int)))))[get_global_id(1)]));\n int j_2 = get_global_id(0);\n int i_3 = get_global_id(1);\n *retval_69 = ((int __global *)(get_region_ptr(mainregion_73, (((int __global *)(get_region_ptr(mainregion_73, (A_1) + (sizeof(int)))))[i_3]) + (sizeof(int)))))[j_2];\n }\n}"));
int main() {
{
region __global * mainregion_73 = create_region(134217728);
int v_71_length_82 = 4;
@eholk
eholk / select-macro.rs
Created August 17, 2012 18:47
Select Macro
{
// select!
macro_rules! select_if {
{
$index:expr,
$count:expr
} => {
fail
../src/test/run-pass/proto-tutorial-code.rs:14:8: 14:19 error: type, variant, or trait must be camel case
../src/test/run-pass/proto-tutorial-code.rs:14 start_state:send { }
^~~~~~~~~~~
../src/test/run-pass/proto-tutorial-code.rs:14:8: 14:19 error: type, variant, or trait must be camel case
../src/test/run-pass/proto-tutorial-code.rs:14 start_state:send { }
^~~~~~~~~~~
../src/test/run-pass/proto-tutorial-code.rs:14:8: 14:19 error: type, variant, or trait must be camel case
../src/test/run-pass/proto-tutorial-code.rs:14 start_state:send { }
^~~~~~~~~~~
../src/test/run-pass/proto-tutorial-code.rs:25:8: 25:19 error: type, variant, or trait must be camel case