Skip to content

Instantly share code, notes, and snippets.

@arsenm
arsenm / memdep_requires_dominator_tree.ll
Created March 30, 2013 02:07
Test case for memdep crash
; ModuleID = 'input.bc'
target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-a0:0:64-f80:32:32"
target triple = "amdil-pc-amdopencl"
@.str = internal addrspace(2) constant [6 x i8] c"uint*\00"
@.str1 = internal addrspace(2) constant [6 x i8] c"uint*\00"
@llvm.argtypename.annotations.__OpenCL_memdep_requires_dominator_tree_kernel = global [2 x i8*] [i8* bitcast ([6 x i8] addrspace(2)* @.str to i8*), i8* bitcast ([6 x i8] addrspace(2)* @.str1 to i8*)], section "llvm.metadata"
@sgv = internal addrspace(2) constant [1 x i8] zeroinitializer
@fgv = internal addrspace(2) constant [1 x i8] zeroinitializer
@lvgv = internal constant [0 x i8*] zeroinitializer
@arsenm
arsenm / as_cv.cpp
Last active December 13, 2015 23:09
Initializer with address_space requires cv qualifier
// No error
__attribute__((address_space(42)))
const float withc = 1.0f;
// No error
__attribute__((address_space(42)))
volatile float withv = 1.0f;
// Error
__attribute__((address_space(42)))
@arsenm
arsenm / as_initializer.c
Created February 19, 2013 21:16
volatile address space initializer
// Compiles successfully
__attribute__((address_space(42)))
volatile float arst[10] = { 0.0f };
/*
as_initializer.cpp:8:20: error: cannot initialize an array element of type
'__attribute__((address_space(42))) float' with an rvalue of type 'float'
float asdf[10] = { 0.0f };
@arsenm
arsenm / CODS scheduler insert
Created August 22, 2012 02:02
Which is better
#if 0
// I'm not sure if it's safe to use the move'd workunit
// after insertion fails, so copy the name so we can say
// which one was problematic
const std::string workunitName = newWorkunit->name();
auto inserted = m_workunits.insert(std::make_pair(newWorkunit->name(), std::move(newWorkunit)));
if (inserted.second) // Insertion succeeded
{
// Actually a new workunit, this is normal. Schedule the inserted workunit
@arsenm
arsenm / gist:1597827
Created January 12, 2012 01:07
VS linker crash
> Finished searching libraries
8> Generating code
8>c:\milkyway_msvc\milkywayathome_client\nbody\src\nbody_plain.c(121): fatal error C1001: An internal error has occurred in the compiler.
8> (compiler file 'f:\dd\vctools\compiler\utc\src\p2\main.c[0x50D245D3:0x000004E6]', line 183)
8> To work around this problem, try simplifying or changing the program near the locations listed above.
8> Please choose the Technical Support command on the Visual C++
8> Help menu, or open the Technical Support help file for more information
8>
8>LINK : fatal error LNK1000: Internal error during IMAGE::BuildImage
8>
@arsenm
arsenm / wrong_size.d
Created January 3, 2012 01:07
Why is the size smaller than it should be?
import std.stdio;
struct Vector(T)
{
T x, y, z;
}
struct Node(T)
{
Vector!(T) position;
@arsenm
arsenm / iterationBug.cl
Created October 18, 2011 16:16
AMD iteration bug?
/* Launch with global size = 6144, local size = 256
Linux x86_64 (kernel 3.0), Catalyst 11.9, SDK 2.5, building for Cayman
*/
#if cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#elif cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#else
@arsenm
arsenm / amd_bug_trio.c
Created October 18, 2011 03:01
Triple AMD compiler bug
/*
3 bugs related to use of an undeclared variable "step"
This does not happen with other names I tried, only step.
Tested: Linux x86_64 (kernel 3.0), Catalyst 11.9, SDK 2.5
Compiling for Cayman
Define BUG to 1, 2, or 3 to experience each bug
*/
@arsenm
arsenm / gist:1252767
Created September 30, 2011 05:29
lol ATISTREAMSDKROOT
/tmp/OCLF8q18i.cl(281): error: invalid argument to attribute
"reqd_work_group_size"
__attribute__ ((reqd_work_group_size(THREADS1, 1, 1)))
^
/tmp/OCLF8q18i.cl(396): error: invalid argument to attribute
"reqd_work_group_size"
__attribute__ ((reqd_work_group_size(THREADS2, 1, 1)))
^
@arsenm
arsenm / CrashTestCase.cl
Created August 11, 2011 23:56
AMD OpenCL compiler crash test case
__kernel void broken(__global double* _x,
__global double* _y,
__global int* debug,
int missing)
{
int k = get_global_id(0);
while (k < 100)
{
//double x = _x[k];