How to create a bitcast instruction?

Hi, I am trying to create a bitcast IR instruction to convert a pointer with a type of float* to void**, Below is the cpp code that I want to use IR to implement:

float * p;
uint64_t size = 1;
hipMalloc((void**)&p,size);

And its corresponding IR is below:

  %22 = alloca float*, align 8            //p is %22
  %23 = alloca float*, align 8
  %24 = alloca i8*, align 8
  %25 = alloca %struct.ihipStream_t*, align 8
  %26 = alloca %struct.ihipEvent_t*, align 8
  %27 = alloca %struct.ihipEvent_t*, align 8
  %28 = alloca float, align 4
  %29 = bitcast float** %22 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %29) #19
  store float* null, float** %22, align 8, !tbaa !23
  %30 = bitcast float** %23 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %30) #19
  store float* null, float** %23, align 8, !tbaa !23
  %31 = bitcast i8** %24 to i8*
  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %31) #19
  store i8* null, i8** %24, align 8, !tbaa !23
  %32 = mul i32 %4, %3
  %33 = add i32 %4, %3
  %34 = mul i32 %33, %5
  %35 = add i32 %34, %32
  %36 = shl i32 %35, 1
  %37 = sext i32 %36 to i64
  %38 = mul nsw i64 %37, 6
  %39 = bitcast float** %22 to i8**                 //Here is the target bitcast inst
  %40 = sext i32 %3 to i64
  %41 = shl nsw i64 %40, 2
  %42 = sext i32 %4 to i64
  %43 = mul i64 %41, %42
  %44 = call i32 @hipMalloc(i8** nonnull %39, i64 %43)           

Now I am stuck where I am trying to create the bitcast IR, I already implement the code to get the i8** pointer type:

                PointerType * int8_ptr_type = Type::getInt8PtrTy(M.getContext());
                if(int8_ptr_type == nullptr)
                {
                    errs()<<"Fail to create int8_ptr_type\n";
                    exit(1);
                }
                PointerType * nested_ptr_type = int8_ptr_type->getPointerTo();
                if(nested_ptr_type == nullptr)
                {
                    errs()<<"Fail to create nested int8** type\n";
                    exit(1);
                }

However, If I just create a bitcast IR, according to its def:

Value *CreateBitCast(Value *V, Type *DestTy,
                       const Twine &Name = "") {
    return CreateCast(Instruction::BitCast, V, DestTy, Name);
  }

I cannot specify the type of input Value to be float**, It may end up generating such a wrong IR instruction: %39 = bitcast float* %22 to i8**. So what should I do to produce the IR to implement the cpp code below?

float * p;
uint64_t size = 1;
hipMalloc((void**)&p,size);

Thank you!

You can’t specify the input type because a Value's type is an intrinsic property, defined by the instruction (or whatever) that creates it.

In this case %22 was created by an alloca float* so its type is already the float** that you want it to be.

Thanks for you answering my question again. I really appreciate your help.
Do you mean that the bitcast instruction can auto-transform the type of %22 from float * to float **? But how can the compiler knows about this, is it changing the input type according to the output type in order to keep them having the same meaning(e.g. nested pointer to nested pointer, pointer to pointer, value to value)?
If so, does it mean that I can directly define a Value based on the float pointer declared before to be the first input arg of hipMalloc(void**,int)?

Not quite. The alloca instruction allocates memory for the type you request (float* here), and returns a pointer to it, i.e. a float**. In C terms the result of the alloca instruction is &p. The compiler isn’t changing anything.

I’m afraid I don’t understand the last question very well. Based on the C you posted, you can directly bitcast %22 from float** to i8** and use it as the argument to hipMalloc. I would expect hipMalloc to store a pointer to the memory it’s allocated there, and you can then do load float*, float** %22 to get the actual float* allocated.

Ok, I see what you mean here. So you are saying the value in %22 is actually the address of p. Well I thought it was the value of p instead of the address of p. That’s my fault, sorry about that.
And yes, about your second paragraph, I already see what you said in my IR code.
Thank you so much! Have a nice day.