cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

kd2
Adept II

ulong16 not implemented fully in libamdocl64??

I've got the following code crashing within the clBuildProgram() function. Specifically gdb SIGSEGVs at

   SCRegSpill::CreateSplitReload(SCInst*, int, unsigned short, SCBlock*, bitset*, bitset*) () from /usr/lib64/libamdocl64.so


I'm using AMD-APP-SDK-v2.8-lnx64.tgz with linux kernel 3.4.4, Catalyst-13.1, on a CL_DEVICE_NAME=Tahiti card (if it matters, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG=1).

I seem to be able to get along just fine coding with ulong8s. It's just when I try to jump to using ulong16s that I eventually run into trouble like this.

Should ulong16s just be avoided for the time being?


-----------------------------------------------------------------------------------------------------------------------

#include <CL/cl.h>

#include <stdio.h>

#include <string.h>

const char *src = "\n\

__kernel void kernel_0(__global ulong16 *in, __global int *out, uint N) \n\

{ \n\

    ulong16 A = in[0];          \n\

    ulong16 B0 = (ulong16)0;    \n\

    ulong16 B1 = (ulong16)0;    \n\

    ulong16 B2 = (ulong16)0;    \n\

    ulong16 B3 = (ulong16)0;    \n\

    for(uint i=0; i < N; i++) { \n\

        B0 += A << 0;           \n\

        B1 += A << 1;           \n\

        B2 += A << 2;           \n\

        B3 += A << 3;           \n\

    }                           \n\

    ulong16 C0 = B0 + B1;       \n\

    ulong16 C1 = B2 + B3;       \n\

    double16 D0 = as_double16(C0);  \n\

    double16 D1 = as_double16(C1);  \n\

    long16 D = isgreater(D0,D1);        \n\

    out[4] = any(D);            \n\

}\n";

int main(int argc, char *argv[])

{

    cl_platform_id platform;

    cl_device_id dev;

    cl_uint platforms, devs;

    clGetPlatformIDs(1,&platform,&platforms);

    clGetDeviceIDs(platform,CL_DEVICE_TYPE_ALL,1,&dev,&devs);

    cl_context_properties properties[] = {CL_CONTEXT_PLATFORM,(cl_context_properties)platform,0};

    cl_int err;

    cl_context ctx = clCreateContext(properties,1,&dev,NULL,NULL,&err);

    size_t src_sz(::strlen(src));

    cl_program prog = clCreateProgramWithSource(ctx,1,&src,&src_sz,&err);

    clBuildProgram(prog,0,NULL,"",NULL,NULL);

}

0 Likes
5 Replies
himanshu_gautam
Grandmaster

Hi kd2,

I could reproduce the crash with 13.1 driver, but clBuildProgram passed without any error when tested on a linux machine with internal driver.

Testcase reproduced on : Tahiti, 13.1 driver, Linux 64-bit (Ubuntu 12.04)

Issue not found on: Capverde GPU, Internal Driver, Linux 64-bit (Ubuntu 12.04)

So hopefully this issue will no longer exist in the next driver. Thanks for reporting it.

0 Likes

Thanks.. but can you test this second kernel also?  The 13.3 beta driver builds the first kernel o.k, and so your internal driver may not be much different. Small tweaks will still crash it.  This following kernel crashes clBuildProgram() in the current 13.3 beta.

const char *src = "\n\

__kernel void kernel_0(__global ulong16 *in, __global int *out, uint N) { \n\

    ulong16 A = in[0],                      \n\

        B0 = (ulong16)0, B1 = (ulong16)0,   \n\

        B2 = (ulong16)0, B3 = (ulong16)0,   \n\

        B4 = (ulong16)0, B5 = (ulong16)0,   \n\

        B6 = (ulong16)0, B7 = (ulong16)0;   \n\

    for(uint i=0; i < N; i++) {             \n\

        B0 += A << 0; B1 += A << 1;         \n\

        B2 += A << 2; B3 += A << 3;         \n\

        B4 += A << 4; B5 += A << 5;         \n\

        B6 += A << 6; B7 += A << 7;         \n\

    }                                       \n\

    out[4] = any(B0+B1+B2+B3==B4+B5+B6+B7); \n\

}\n";

0 Likes

The function any as described in section 6.12.6 in OpenCL spec 1.2 only takes igenType as its argument( which can be char, short, int, long and their vector counterparts). I did a typecast of the argument to int, and the kernel does not crash.

Nevertheless, OCL compiler should not crash, and give relevant error in this case. I will forward it to AMD Engg Team.

0 Likes

No, the code is according to spec. The == operator (OpenCL spec 6.3e) take source vector operands of type ulongn and returns a longn.  The any() function (6.12.6) should take that operand of longn.

0 Likes

Yeah right. Thanks for the information. It was useful.

0 Likes